From a518e7063df64ed3564ef9bf5dcfda6702d6d54b Mon Sep 17 00:00:00 2001 From: Dmitry Budnikov Date: Fri, 23 Nov 2018 17:51:15 +0300 Subject: [PATCH] Merge pull request #13120 from dbudniko:dbudniko/gapi_opencl_kernel_example * custom OpenCL G-API kernel draft * clean up and warnings fix * more warnings * white space * new blank line at the EOF removed * HAVE_OPENCL guard * remove unnecessary ocl API call * remove sum test workaround * check if opencl activated * fix std::str warning * CPU fall back for symm7x7 * gpu test kernel draft * adjust have opencl guard * more guards * one more attempt to adjust guards * empty stub files and kernel source files creation in the test directory * try to force auto generation * one more attempt to force build * remove symm7x7 custom from gapi module * looks like that this version works properly on Win desktop * clean up * more clean up * address some suggestions from Dmitry's review * const kernel coefficients * CV_Error in kernel + try to fix cpu fallback * CV_Error_ instead CV_Error * everything in one gapi_gpu_test.cpp * fix warning * remove kernel generation, add kernel string * avoid generated code and ocl internal namespace * fix misprint * c_str --- modules/gapi/test/common/gapi_core_tests_inl.hpp | 3 +- modules/gapi/test/gapi_gpu_test.cpp | 207 ++++++++++++++++++ modules/gapi/test/opencl_kernels_test_gapi.hpp | 260 +++++++++++++++++++++++ 3 files changed, 468 insertions(+), 2 deletions(-) create mode 100644 modules/gapi/test/gapi_gpu_test.cpp create mode 100644 modules/gapi/test/opencl_kernels_test_gapi.hpp diff --git a/modules/gapi/test/common/gapi_core_tests_inl.hpp b/modules/gapi/test/common/gapi_core_tests_inl.hpp index d33b5cc..1f07131 100644 --- a/modules/gapi/test/common/gapi_core_tests_inl.hpp +++ b/modules/gapi/test/common/gapi_core_tests_inl.hpp @@ -684,8 +684,7 @@ TEST_P(SumTest, AccuracyTest) cv::Size sz_in = std::get<1>(param); auto tolerance = std::get<3>(param); auto compile_args = std::get<4>(param); - //initMatrixRandU(std::get<0>(param), sz_in, std::get<2>(param)); - initMatsRandN(std::get<0>(param), sz_in, std::get<2>(param)); //TODO: workaround trying to fix SumTest failures + initMatrixRandU(std::get<0>(param), sz_in, std::get<2>(param)); cv::Scalar out_sum; diff --git a/modules/gapi/test/gapi_gpu_test.cpp b/modules/gapi/test/gapi_gpu_test.cpp new file mode 100644 index 0000000..7cb6f9f --- /dev/null +++ b/modules/gapi/test/gapi_gpu_test.cpp @@ -0,0 +1,207 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. +// +// Copyright (C) 2018 Intel Corporation + + +#include "test_precomp.hpp" + + +#include "logger.hpp" +#include "common/gapi_tests_common.hpp" +#include "opencv2/gapi/gpu/ggpukernel.hpp" +#include "opencl_kernels_test_gapi.hpp" + + +namespace cv +{ + +#ifdef HAVE_OPENCL + + static void reference_symm7x7_CPU(const cv::Mat& in, const cv::Mat& kernel_coeff, int shift, cv::Mat &out) + { + cv::Point anchor = { -1, -1 }; + double delta = 0; + + const int* ci = kernel_coeff.ptr(); + + float c_float[10]; + float divisor = (float)(1 << shift); + for (int i = 0; i < 10; i++) + { + c_float[i] = ci[i] / divisor; + } + // J & I & H & G & H & I & J + // I & F & E & D & E & F & I + // H & E & C & B & C & E & H + // G & D & B & A & B & D & G + // H & E & C & B & C & E & H + // I & F & E & D & E & F & I + // J & I & H & G & H & I & J + + // A & B & C & D & E & F & G & H & I & J + + // 9 & 8 & 7 & 6 & 7 & 8 & 9 + // 8 & 5 & 4 & 3 & 4 & 5 & 8 + // 7 & 4 & 2 & 1 & 2 & 4 & 7 + // 6 & 3 & 1 & 0 & 1 & 3 & 6 + // 7 & 4 & 2 & 1 & 2 & 4 & 7 + // 8 & 5 & 4 & 3 & 4 & 5 & 8 + // 9 & 8 & 7 & 6 & 7 & 8 & 9 + + float coefficients[49] = + { + c_float[9], c_float[8], c_float[7], c_float[6], c_float[7], c_float[8], c_float[9], + c_float[8], c_float[5], c_float[4], c_float[3], c_float[4], c_float[5], c_float[8], + c_float[7], c_float[4], c_float[2], c_float[1], c_float[2], c_float[4], c_float[7], + c_float[6], c_float[3], c_float[1], c_float[0], c_float[1], c_float[3], c_float[6], + c_float[7], c_float[4], c_float[2], c_float[1], c_float[2], c_float[4], c_float[7], + c_float[8], c_float[5], c_float[4], c_float[3], c_float[4], c_float[5], c_float[8], + c_float[9], c_float[8], c_float[7], c_float[6], c_float[7], c_float[8], c_float[9] + }; + + cv::Mat kernel = cv::Mat(7, 7, CV_32FC1); + float* cf = kernel.ptr(); + for (int i = 0; i < 49; i++) + { + cf[i] = coefficients[i]; + } + + cv::filter2D(in, out, CV_8UC1, kernel, anchor, delta, cv::BORDER_REPLICATE); + } + + namespace gapi_test_kernels + { + G_TYPED_KERNEL(TSymm7x7_test, , "org.opencv.imgproc.symm7x7_test") { + static GMatDesc outMeta(GMatDesc in, Mat, int) { + return in.withType(CV_8U, 1); + } + }; + + + GAPI_GPU_KERNEL(GGPUSymm7x7_test, TSymm7x7_test) + { + static void run(const cv::UMat& in, const cv::Mat& kernel_coeff, int shift, cv::UMat &out) + { + if (cv::ocl::isOpenCLActivated()) + { + cv::Size size = in.size(); + size_t globalsize[2] = { (size_t)size.width, (size_t)size.height }; + + const cv::String moduleName = "gapi"; + cv::ocl::ProgramSource source(moduleName, "symm7x7", opencl_symm7x7_src, ""); + + static const char * const borderMap[] = { "BORDER_CONSTANT", "BORDER_REPLICATE", "BORDER_UNDEFINED" }; + std::string build_options = " -D BORDER_CONSTANT_VALUE=" + std::to_string(0) + + " -D " + borderMap[1] + + " -D SCALE=1.f/" + std::to_string(1 << shift) + ".f"; + + cv::String errmsg; + cv::ocl::Program program(source, build_options, errmsg); + if (program.ptr() == NULL) + { + CV_Error_(cv::Error::OpenCLInitError, ("symm_7x7_test Can't compile OpenCL program: = %s with build_options = %s\n", errmsg.c_str(), build_options.c_str())); + } + if (!errmsg.empty()) + { + std::cout << "OpenCL program build log:" << std::endl << errmsg << std::endl; + } + + cv::ocl::Kernel kernel("symm_7x7_test", program); + if (kernel.empty()) + { + CV_Error(cv::Error::OpenCLInitError, "symm_7x7_test Can't get OpenCL kernel\n"); + } + + cv::UMat gKer; + kernel_coeff.copyTo(gKer); + + int tile_y = 0; + + int idxArg = kernel.set(0, cv::ocl::KernelArg::PtrReadOnly(in)); + idxArg = kernel.set(idxArg, (int)in.step); + idxArg = kernel.set(idxArg, (int)size.width); + idxArg = kernel.set(idxArg, (int)size.height); + idxArg = kernel.set(idxArg, cv::ocl::KernelArg::PtrWriteOnly(out)); + idxArg = kernel.set(idxArg, (int)out.step); + idxArg = kernel.set(idxArg, (int)size.height); + idxArg = kernel.set(idxArg, (int)size.width); + idxArg = kernel.set(idxArg, (int)tile_y); + idxArg = kernel.set(idxArg, cv::ocl::KernelArg::PtrReadOnly(gKer)); + + if (!kernel.run(2, globalsize, NULL, false)) + { + CV_Error(cv::Error::OpenCLApiCallError, "symm_7x7_test OpenCL kernel run failed\n"); + } + } + else + { + //CPU fallback + cv::Mat in_Mat, out_Mat; + in_Mat = in.getMat(ACCESS_READ); + out_Mat = out.getMat(ACCESS_WRITE); + reference_symm7x7_CPU(in_Mat, kernel_coeff, shift, out_Mat); + } + } + }; + + cv::gapi::GKernelPackage gpuTestPackage = cv::gapi::kernels + (); + + } // namespace gapi_test_kernels +#endif //HAVE_OPENCL + +} // namespace cv + + +namespace opencv_test +{ + +#ifdef HAVE_OPENCL + +using namespace cv::gapi_test_kernels; + +TEST(GPU, Symm7x7_test) +{ + const auto sz = cv::Size(1280, 720); + cv::Mat in_mat = cv::Mat::eye(sz, CV_8UC1); + cv::Mat out_mat_gapi(sz, CV_8UC1); + cv::Mat out_mat_ocv(sz, CV_8UC1); + cv::Scalar mean = cv::Scalar(127.0f); + cv::Scalar stddev = cv::Scalar(40.f); + cv::randn(in_mat, mean, stddev); + + //Symm7x7 coefficients and shift + int coefficients_symm7x7[10] = { 1140, -118, 526, 290, -236, 64, -128, -5, -87, -7 }; + int shift = 10; + cv::Mat kernel_coeff(10, 1, CV_32S); + int* ci = kernel_coeff.ptr(); + for (int i = 0; i < 10; i++) + { + ci[i] = coefficients_symm7x7[i]; + } + + // Run G-API + cv::GMat in; + auto out = TSymm7x7_test::on(in, kernel_coeff, shift); + cv::GComputation comp(cv::GIn(in), cv::GOut(out)); + + auto cc = comp.compile(cv::descr_of(in_mat), cv::compile_args(gpuTestPackage)); + cc(cv::gin(in_mat), cv::gout(out_mat_gapi)); + + // Run OpenCV + reference_symm7x7_CPU(in_mat, kernel_coeff, shift, out_mat_ocv); + + compare_f cmpF = AbsSimilarPoints(1, 0.05).to_compare_f(); + + // Comparison ////////////////////////////////////////////////////////////// + { + EXPECT_TRUE(cmpF(out_mat_gapi, out_mat_ocv)); + EXPECT_EQ(out_mat_gapi.size(), sz); + } +} +#endif + +} // namespace opencv_test diff --git a/modules/gapi/test/opencl_kernels_test_gapi.hpp b/modules/gapi/test/opencl_kernels_test_gapi.hpp new file mode 100644 index 0000000..1164165 --- /dev/null +++ b/modules/gapi/test/opencl_kernels_test_gapi.hpp @@ -0,0 +1,260 @@ +// This file is part of OpenCV project. +// It is subject to the license terms in the LICENSE file found in the top-level directory +// of this distribution and at http://opencv.org/license.html. +// +// Copyright (C) 2018 Intel Corporation + +#include "opencv2/core/ocl.hpp" +#include "opencv2/core/ocl_genbase.hpp" +#include "opencv2/core/opencl/ocl_defs.hpp" + +#ifdef HAVE_OPENCL +const char* opencl_symm7x7_src = +"#if BORDER_REPLICATE\n" +"#define GET_BORDER(elem) (elem)\n" +"#define SET_ALL(i, j) a0[i] = a0[j]; a1[i] = a1[j]; a2[i] = a2[j]; b[i] = b[j]; c0[i] = c0[j]; c1[i] = c1[j]; c2[i] = c2[j];\n" +"#else\n" +"#define GET_BORDER(elem) (BORDER_CONSTANT_VALUE)\n" +"#define SET_ALL(i, j) a0[i] = a1[i] = a2[i] = c0[i] = c1[i] = c2[i] = BORDER_CONSTANT_VALUE; b[i] = BORDER_CONSTANT_VALUE;\n" +"#endif\n" +"#define GET_A0(id, x, l_edge, a1) ((x) <= (l_edge + 2) ? GET_BORDER(a1) : (((const __global uchar*)(id))[-3]))\n" +"#define GET_A1(id, x, l_edge, a2) ((x) <= (l_edge + 1) ? GET_BORDER(a2) : (((const __global uchar*)(id))[-2]))\n" +"#define GET_A2(id, x, l_edge, b) ((x) <= (l_edge) ? GET_BORDER(b[0]) : (((const __global uchar*)(id))[-1]))\n" +"#define GET_C0(id, x, r_edge, b) ((x) >= (r_edge) ? GET_BORDER(b[8 - 1]) : (((const __global uchar*)(id))[8]))\n" +"#define GET_C1(id, x, r_edge, c0) ((x) >= (r_edge - 1) ? GET_BORDER(c0) : (((const __global uchar*)(id))[8 + 1]))\n" +"#define GET_C2(id, x, r_edge, c1) ((x) >= (r_edge - 2) ? GET_BORDER(c1) : (((const __global uchar*)(id))[8 + 2]))\n" +"__kernel void symm_7x7_test(\n" +"__global const uchar * srcptr,\n" +"int srcStep, int srcEndX, int srcEndY,\n" +"__global uchar * dstptr, int dstStep,\n" +"int rows, int cols,\n" +"int tile_y_coord,\n" +"__constant int * coeff)\n" +"{\n" +"int lEdge = 0, rEdge = cols - 8;\n" +"int x = (get_global_id(0) < cols/8) ? get_global_id(0) * 8: cols - 8;\n" +"int y = get_global_id(1);\n" +"int yd = min(3, tile_y_coord);\n" +"int dst_id = mad24(y, dstStep, x);\n" +"y+=yd;\n" +"int src_id = mad24(y, srcStep, x);\n" +"int y_limit = y + tile_y_coord;\n" +"y_limit-=yd;\n" +"const __global uchar* psrc = (const __global uchar*)(srcptr + src_id);\n" +"__global uchar* pdst = (__global uchar*)(dstptr + dst_id);\n" +"#define BSIZE (7)\n" +"float a0[BSIZE]; float a1[BSIZE]; float a2[BSIZE];\n" +"float8 b[BSIZE];\n" +"float c0[BSIZE]; float c1[BSIZE]; float c2[BSIZE];\n" +"b[3] = convert_float8(vload8(0, (const __global uchar*)psrc));\n" +"if( (y_limit <=2 ) || (y_limit >= srcEndY - 3) || (x >= rEdge-2) || (x <= lEdge + 2) )\n" +"{\n" +"a2[3] = GET_A2(psrc, x, lEdge, b[3]);\n" +"a1[3] = GET_A1(psrc, x, lEdge, a2[3]);\n" +"a0[3] = GET_A0(psrc, x, lEdge, a1[3]);\n" +"c0[3] = GET_C0(psrc, x, rEdge, b[3]);\n" +"c1[3] = GET_C1(psrc, x, rEdge, c0[3]);\n" +"c2[3] = GET_C2(psrc, x, rEdge, c1[3]);\n" +"if(y_limit > 0)\n" +"{\n" +"b[2] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep)));\n" +"a2[2] = GET_A2(psrc - srcStep, x, lEdge, b[2]);\n" +"a1[2] = GET_A1(psrc - srcStep, x, lEdge, a2[2]);\n" +"a0[2] = GET_A0(psrc - srcStep, x, lEdge, a1[2]);\n" +"c0[2] = GET_C0(psrc - srcStep, x, rEdge, b[2]);\n" +"c1[2] = GET_C1(psrc - srcStep, x, rEdge, c0[2]);\n" +"c2[2] = GET_C2(psrc - srcStep, x, rEdge, c1[2]);\n" +"}\n" +"else\n" +"{\n" +"SET_ALL(2, 3);\n" +"}\n" +"if( y_limit > 1 )\n" +"{\n" +"b[1] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*2)));\n" +"a2[1] = GET_A2(psrc - srcStep*2, x, lEdge, b[1]);\n" +"a1[1] = GET_A1(psrc - srcStep*2, x, lEdge, a2[1]);\n" +"a0[1] = GET_A0(psrc - srcStep*2, x, lEdge, a1[1]);\n" +"c0[1] = GET_C0(psrc - srcStep*2, x, rEdge, b[1]);\n" +"c1[1] = GET_C1(psrc - srcStep*2, x, rEdge, c0[1]);\n" +"c2[1] = GET_C2(psrc - srcStep*2, x, rEdge, c1[1]);\n" +"}\n" +"else\n" +"{\n" +"SET_ALL(1, 2);\n" +"}\n" +"if( y_limit > 2 )\n" +"{\n" +"b[0] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*3)));\n" +"a2[0] = GET_A2(psrc - srcStep*3, x, lEdge, b[0]);\n" +"a1[0] = GET_A1(psrc - srcStep*3, x, lEdge, a2[0]);\n" +"a0[0] = GET_A0(psrc - srcStep*3, x, lEdge, a1[0]);\n" +"c0[0] = GET_C0(psrc - srcStep*3, x, rEdge, b[0]);\n" +"c1[0] = GET_C1(psrc - srcStep*3, x, rEdge, c0[0]);\n" +"c2[0] = GET_C2(psrc - srcStep*3, x, rEdge, c1[0]);\n" +"}\n" +"else\n" +"{\n" +"SET_ALL(0, 1);\n" +"}\n" +"if( y_limit < srcEndY - 1 )\n" +"{\n" +"b[4] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep)));\n" +"a2[4] = GET_A2(psrc + srcStep, x, lEdge, b[4]);\n" +"a1[4] = GET_A1(psrc + srcStep, x, lEdge, a2[4]);\n" +"a0[4] = GET_A0(psrc + srcStep, x, lEdge, a1[4]);\n" +"c0[4] = GET_C0(psrc + srcStep, x, rEdge, b[4]);\n" +"c1[4] = GET_C1(psrc + srcStep, x, rEdge, c0[4]);\n" +"c2[4] = GET_C2(psrc + srcStep, x, rEdge, c1[4]);\n" +"}\n" +"else\n" +"{\n" +"SET_ALL(4, 3);\n" +"}\n" +"if( y_limit < srcEndY - 2 )\n" +"{\n" +"b[5] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*2)));\n" +"a2[5] = GET_A2(psrc + srcStep*2, x, lEdge, b[5]);\n" +"a1[5] = GET_A1(psrc + srcStep*2, x, lEdge, a2[5]);\n" +"a0[5] = GET_A0(psrc + srcStep*2, x, lEdge, a1[5]);\n" +"c0[5] = GET_C0(psrc + srcStep*2, x, rEdge, b[5]);\n" +"c1[5] = GET_C1(psrc + srcStep*2, x, rEdge, c0[5]);\n" +"c2[5] = GET_C2(psrc + srcStep*2, x, rEdge, c1[5]);\n" +"}\n" +"else\n" +"{\n" +"SET_ALL(5, 4);\n" +"}\n" +"if( y_limit < srcEndY - 3 )\n" +"{\n" +"b[6] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*3)));\n" +"a2[6] = GET_A2(psrc + srcStep*3, x, lEdge, b[6]);\n" +"a1[6] = GET_A1(psrc + srcStep*3, x, lEdge, a2[6]);\n" +"a0[6] = GET_A0(psrc + srcStep*3, x, lEdge, a1[6]);\n" +"c0[6] = GET_C0(psrc + srcStep*3, x, rEdge, b[6]);\n" +"c1[6] = GET_C1(psrc + srcStep*3, x, rEdge, c0[6]);\n" +"c2[6] = GET_C2(psrc + srcStep*3, x, rEdge, c1[6]);\n" +"}\n" +"else\n" +"{\n" +"SET_ALL(6, 5);\n" +"}\n" +"}\n" +"else\n" +"{\n" +"a2[3] = (((const __global uchar*)(psrc))[-1]);\n" +"a1[3] = (((const __global uchar*)(psrc))[-2]);\n" +"a0[3] = (((const __global uchar*)(psrc))[-3]);\n" +"c0[3] = (((const __global uchar*)(psrc))[8]);\n" +"c1[3] = (((const __global uchar*)(psrc))[8 + 1]);\n" +"c2[3] = (((const __global uchar*)(psrc))[8 + 2]);\n" +"b[2] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep)));\n" +"a2[2] = (((const __global uchar*)(psrc - srcStep))[-1]);\n" +"a1[2] = (((const __global uchar*)(psrc - srcStep))[-2]);\n" +"a0[2] = (((const __global uchar*)(psrc - srcStep))[-3]);\n" +"c0[2] = (((const __global uchar*)(psrc - srcStep))[8]);\n" +"c1[2] = (((const __global uchar*)(psrc - srcStep))[8 + 1]);\n" +"c2[2] = (((const __global uchar*)(psrc - srcStep))[8 + 2]);\n" +"b[1] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*2)));\n" +"a2[1] = (((const __global uchar*)(psrc - srcStep*2))[-1]);\n" +"a1[1] = (((const __global uchar*)(psrc - srcStep*2))[-2]);\n" +"a0[1] = (((const __global uchar*)(psrc - srcStep*2))[-3]);\n" +"c0[1] = (((const __global uchar*)(psrc - srcStep*2))[8]);\n" +"c1[1] = (((const __global uchar*)(psrc - srcStep*2))[8 + 1]);\n" +"c2[1] = (((const __global uchar*)(psrc - srcStep*2))[8 + 2]);\n" +"b[0] = convert_float8(vload8(0, (const __global uchar*)(psrc - srcStep*3)));\n" +"a2[0] = (((const __global uchar*)(psrc - srcStep*3))[-1]);\n" +"a1[0] = (((const __global uchar*)(psrc - srcStep*3))[-2]);\n" +"a0[0] = (((const __global uchar*)(psrc - srcStep*3))[-3]);\n" +"c0[0] = (((const __global uchar*)(psrc - srcStep*3))[8]);\n" +"c1[0] = (((const __global uchar*)(psrc - srcStep*3))[8 + 1]);\n" +"c2[0] = (((const __global uchar*)(psrc - srcStep*3))[8 + 2]);\n" +"b[4] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep)));\n" +"a2[4] = (((const __global uchar*)(psrc + srcStep))[-1]);\n" +"a1[4] = (((const __global uchar*)(psrc + srcStep))[-2]);\n" +"a0[4] = (((const __global uchar*)(psrc + srcStep))[-3]);\n" +"c0[4] = (((const __global uchar*)(psrc + srcStep))[8]);\n" +"c1[4] = (((const __global uchar*)(psrc + srcStep))[8 + 1]);\n" +"c2[4] = (((const __global uchar*)(psrc + srcStep))[8 + 2]);\n" +"b[5] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*2)));\n" +"a2[5] = (((const __global uchar*)(psrc + srcStep*2))[-1]);\n" +"a1[5] = (((const __global uchar*)(psrc + srcStep*2))[-2]);\n" +"a0[5] = (((const __global uchar*)(psrc + srcStep*2))[-3]);\n" +"c0[5] = (((const __global uchar*)(psrc + srcStep*2))[8]);\n" +"c1[5] = (((const __global uchar*)(psrc + srcStep*2))[8 + 1]);\n" +"c2[5] = (((const __global uchar*)(psrc + srcStep*2))[8 + 2]);\n" +"b[6] = convert_float8(vload8(0, (const __global uchar*)(psrc + srcStep*3)));\n" +"a2[6] = (((const __global uchar*)(psrc + srcStep*3))[-1]);\n" +"a1[6] = (((const __global uchar*)(psrc + srcStep*3))[-2]);\n" +"a0[6] = (((const __global uchar*)(psrc + srcStep*3))[-3]);\n" +"c0[6] = (((const __global uchar*)(psrc + srcStep*3))[8]);\n" +"c1[6] = (((const __global uchar*)(psrc + srcStep*3))[8 + 1]);\n" +"c2[6] = (((const __global uchar*)(psrc + srcStep*3))[8 + 2]);\n" +"}\n" +"float a0_sum[3]; float a1_sum[3]; float a2_sum[3];\n" +"float8 b_sum[3];\n" +"float c0_sum[3]; float c1_sum[3]; float c2_sum[3];\n" +"a0_sum[0] = a0[0] + a0[6];\n" +"a0_sum[1] = a0[1] + a0[5];\n" +"a0_sum[2] = a0[2] + a0[4];\n" +"a1_sum[0] = a1[0] + a1[6];\n" +"a1_sum[1] = a1[1] + a1[5];\n" +"a1_sum[2] = a1[2] + a1[4];\n" +"a2_sum[0] = a2[0] + a2[6];\n" +"a2_sum[1] = a2[1] + a2[5];\n" +"a2_sum[2] = a2[2] + a2[4];\n" +"c0_sum[0] = c0[0] + c0[6];\n" +"c0_sum[1] = c0[1] + c0[5];\n" +"c0_sum[2] = c0[2] + c0[4];\n" +"c1_sum[0] = c1[0] + c1[6];\n" +"c1_sum[1] = c1[1] + c1[5];\n" +"c1_sum[2] = c1[2] + c1[4];\n" +"c2_sum[0] = c2[0] + c2[6];\n" +"c2_sum[1] = c2[1] + c2[5];\n" +"c2_sum[2] = c2[2] + c2[4];\n" +"b_sum[0] = b[0] + b[6];\n" +"b_sum[1] = b[1] + b[5];\n" +"b_sum[2] = b[2] + b[4];\n" +"float8 A = b[3];\n" +"float8 intermediate = A * (float)coeff[0];\n" +"float8 B = b_sum[2] +\n" +"(float8)(a2[3], b[3].s0123, b[3].s456) +\n" +"(float8)(b[3].s123, b[3].s4567, c0[3]);\n" +"intermediate += B * (float)coeff[1];\n" +"float8 C = (float8)(a2_sum[2], b_sum[2].s0123, b_sum[2].s456) +\n" +"(float8)(b_sum[2].s123, b_sum[2].s4567, c0_sum[2]);\n" +"intermediate += C * (float)coeff[2];\n" +"float8 D = b_sum[1] +\n" +"(float8)(a1[3], a2[3], b[3].s0123, b[3].s45) +\n" +"(float8)(b[3].s23, b[3].s4567, c0[3], c1[3]);\n" +"intermediate += D * (float)coeff[3];\n" +"float8 E = (float8)(a2_sum[1], b_sum[1].s0123, b_sum[1].s456) +\n" +"(float8)( b_sum[1].s123, b_sum[1].s4567, c0_sum[1]) +\n" +"(float8)( a1_sum[2], a2_sum[2], b_sum[2].s0123, b_sum[2].s45) +\n" +"(float8)( b_sum[2].s23, b_sum[2].s4567, c0_sum[2], c1_sum[2]);\n" +"intermediate += E * (float)coeff[4];\n" +"float8 F = (float8)(a1_sum[1], a2_sum[1], b_sum[1].s0123, b_sum[1].s45) +\n" +"(float8)(b_sum[1].s23, b_sum[1].s4567, c0_sum[1], c1_sum[1]);\n" +"intermediate += F * (float)coeff[5];\n" +"float8 G = b_sum[0] +\n" +"(float8)(a0[3], a1[3], a2[3], b[3].s0123, b[3].s4) +\n" +"(float8)(b[3].s3, b[3].s4567, c0[3], c1[3], c2[3]);\n" +"intermediate += G * (float)coeff[6];\n" +"float8 H = (float8)(a2_sum[0], b_sum[0].s0123, b_sum[0].s456) +\n" +"(float8)(b_sum[0].s123, b_sum[0].s4567, c0_sum[0]) +\n" +"(float8)(a0_sum[2], a1_sum[2], a2_sum[2], b_sum[2].s0123, b_sum[2].s4) +\n" +"(float8)(b_sum[2].s3, b_sum[2].s4567, c0_sum[2], c1_sum[2], c2_sum[2]);\n" +"intermediate += H * (float)coeff[7];\n" +"float8 I = (float8)(a1_sum[0], a2_sum[0], b_sum[0].s0123, b_sum[0].s45) +\n" +"(float8)(b_sum[0].s23, b_sum[0].s4567, c0_sum[0], c1_sum[0]) +\n" +"(float8)(a0_sum[1], a1_sum[1], a2_sum[1], b_sum[1].s0123, b_sum[1].s4) +\n" +"(float8)(b_sum[1].s3, b_sum[1].s4567, c0_sum[1], c1_sum[1], c2_sum[1]);\n" +"intermediate += I * (float)coeff[8];\n" +"float8 J = (float8)(a0_sum[0], a1_sum[0], a2_sum[0], b_sum[0].s0123, b_sum[0].s4) +\n" +"(float8)(b_sum[0].s3, b_sum[0].s4567, c0_sum[0], c1_sum[0], c2_sum[0]);\n" +"intermediate += J * (float)coeff[9];\n" +"intermediate *= SCALE;\n" +"vstore8(convert_uchar8_sat(intermediate), 0, (__global uchar*)(pdst));\n" +"}\n" +; +#endif -- 2.7.4