1 /*M///////////////////////////////////////////////////////////////////////////////////////
3 // IMPORTANT: READ BEFORE DOWNLOADING, COPYING, INSTALLING OR USING.
5 // By downloading, copying, installing or using the software you agree to this license.
6 // If you do not agree to this license, do not download, install,
7 // copy or use the software.
11 // For Open Source Computer Vision Library
13 // Copyright (C) 2010-2012, Multicoreware, Inc., all rights reserved.
14 // Copyright (C) 2010-2012, Advanced Micro Devices, Inc., all rights reserved.
15 // Third party copyrights are property of their respective owners.
18 // Peng Xiao, pengxiao@multicorewareinc.com
20 // Redistribution and use in source and binary forms, with or without modification,
21 // are permitted provided that the following conditions are met:
23 // * Redistribution's of source code must retain the above copyright notice,
24 // this list of conditions and the following disclaimer.
26 // * Redistribution's in binary form must reproduce the above copyright notice,
27 // this list of conditions and the following disclaimer in the documentation
28 // and/or other oclMaterials provided with the distribution.
30 // * The name of the copyright holders may not be used to endorse or promote products
31 // derived from this software without specific prior written permission.
33 // This software is provided by the copyright holders and contributors as is and
34 // any express or implied warranties, including, but not limited to, the implied
35 // warranties of merchantability and fitness for a particular purpose are disclaimed.
36 // In no event shall the Intel Corporation or contributors be liable for any direct,
37 // indirect, incidental, special, exemplary, or consequential damages
38 // (including, but not limited to, procurement of substitute goods or services;
39 // loss of use, data, or profits; or business interruption) however caused
40 // and on any theory of liability, whether in contract, strict liability,
41 // or tort (including negligence or otherwise) arising in any way out of
42 // the use of this software, even if advised of the possibility of such damage.
46 #include "precomp.hpp"
49 #if ((__GNUC__ * 100) + __GNUC_MINOR__) >= 402
50 #define GCC_DIAG_STR(s) #s
51 #define GCC_DIAG_JOINSTR(x,y) GCC_DIAG_STR(x ## y)
52 # define GCC_DIAG_DO_PRAGMA(x) _Pragma (#x)
53 # define GCC_DIAG_PRAGMA(x) GCC_DIAG_DO_PRAGMA(GCC diagnostic x)
54 # if ((__GNUC__ * 100) + __GNUC_MINOR__) >= 406
55 # define GCC_DIAG_OFF(x) GCC_DIAG_PRAGMA(push) \
56 GCC_DIAG_PRAGMA(ignored GCC_DIAG_JOINSTR(-W,x))
57 # define GCC_DIAG_ON(x) GCC_DIAG_PRAGMA(pop)
59 # define GCC_DIAG_OFF(x) GCC_DIAG_PRAGMA(ignored GCC_DIAG_JOINSTR(-W,x))
60 # define GCC_DIAG_ON(x) GCC_DIAG_PRAGMA(warning GCC_DIAG_JOINSTR(-W,x))
63 # define GCC_DIAG_OFF(x)
64 # define GCC_DIAG_ON(x)
75 inline int divUp(int total, int grain)
77 return (total + grain - 1) / grain;
80 // provide additional methods for the user to interact with the command queue after a task is fired
81 static void openCLExecuteKernel_2(Context *clCxt , const char **source, string kernelName, size_t globalThreads[3],
82 size_t localThreads[3], vector< pair<size_t, const void *> > &args, int channels,
83 int depth, char *build_options, FLUSH_MODE finish_mode)
85 //construct kernel name
86 //The rule is functionName_Cn_Dn, C represent Channels, D Represent DataType Depth, n represent an integer number
87 //for exmaple split_C2_D2, represent the split kernel with channels =2 and dataType Depth = 2(Data type is char)
90 idxStr << "_C" << channels;
92 idxStr << "_D" << depth;
93 kernelName += idxStr.str();
96 kernel = openCLGetKernelFromSource(clCxt, source, kernelName, build_options);
98 if ( localThreads != NULL)
100 globalThreads[0] = divUp(globalThreads[0], localThreads[0]) * localThreads[0];
101 globalThreads[1] = divUp(globalThreads[1], localThreads[1]) * localThreads[1];
102 globalThreads[2] = divUp(globalThreads[2], localThreads[2]) * localThreads[2];
104 //size_t blockSize = localThreads[0] * localThreads[1] * localThreads[2];
105 cv::ocl::openCLVerifyKernel(clCxt, kernel, localThreads);
107 for(size_t i = 0; i < args.size(); i ++)
108 openCLSafeCall(clSetKernelArg(kernel, i, args[i].first, args[i].second));
110 openCLSafeCall(clEnqueueNDRangeKernel((cl_command_queue)clCxt->oclCommandQueue(), kernel, 3, NULL, globalThreads,
111 localThreads, 0, NULL, NULL));
116 clFinish((cl_command_queue)clCxt->oclCommandQueue());
118 clFlush((cl_command_queue)clCxt->oclCommandQueue());
124 openCLSafeCall(clReleaseKernel(kernel));
127 void openCLExecuteKernel2(Context *clCxt , const char **source, string kernelName,
128 size_t globalThreads[3], size_t localThreads[3],
129 vector< pair<size_t, const void *> > &args, int channels, int depth, FLUSH_MODE finish_mode)
131 openCLExecuteKernel2(clCxt, source, kernelName, globalThreads, localThreads, args,
132 channels, depth, NULL, finish_mode);
134 void openCLExecuteKernel2(Context *clCxt , const char **source, string kernelName,
135 size_t globalThreads[3], size_t localThreads[3],
136 vector< pair<size_t, const void *> > &args, int channels, int depth, char *build_options, FLUSH_MODE finish_mode)
139 openCLExecuteKernel_2(clCxt, source, kernelName, globalThreads, localThreads, args, channels, depth,
140 build_options, finish_mode);
144 GCC_DIAG_OFF(deprecated-declarations)
146 cl_mem bindTexture(const oclMat &mat)
149 cl_image_format format;
151 int depth = mat.depth();
152 int channels = mat.channels();
157 format.image_channel_data_type = CL_UNSIGNED_INT8;
160 format.image_channel_data_type = CL_UNSIGNED_INT32;
163 format.image_channel_data_type = CL_FLOAT;
166 CV_Error(-1, "Image forma is not supported");
172 format.image_channel_order = CL_R;
175 format.image_channel_order = CL_RGB;
178 format.image_channel_order = CL_RGBA;
181 CV_Error(-1, "Image format is not supported");
184 #ifdef CL_VERSION_1_2
185 //this enables backwards portability to
186 //run on OpenCL 1.1 platform if library binaries are compiled with OpenCL 1.2 support
187 if(Context::getContext()->supportsFeature(Context::CL_VER_1_2))
190 desc.image_type = CL_MEM_OBJECT_IMAGE2D;
191 desc.image_width = mat.cols;
192 desc.image_height = mat.rows;
193 desc.image_depth = 0;
194 desc.image_array_size = 1;
195 desc.image_row_pitch = 0;
196 desc.image_slice_pitch = 0;
198 desc.num_mip_levels = 0;
199 desc.num_samples = 0;
200 texture = clCreateImage((cl_context)mat.clCxt->oclContext(), CL_MEM_READ_WRITE, &format, &desc, NULL, &err);
205 texture = clCreateImage2D(
206 (cl_context)mat.clCxt->oclContext(),
215 size_t origin[] = { 0, 0, 0 };
216 size_t region[] = { mat.cols, mat.rows, 1 };
219 if (mat.cols * mat.elemSize() != mat.step)
221 devData = clCreateBuffer((cl_context)mat.clCxt->oclContext(), CL_MEM_READ_ONLY, mat.cols * mat.rows
222 * mat.elemSize(), NULL, NULL);
223 const size_t regin[3] = {mat.cols * mat.elemSize(), mat.rows, 1};
224 clEnqueueCopyBufferRect((cl_command_queue)mat.clCxt->oclCommandQueue(), (cl_mem)mat.data, devData, origin, origin,
225 regin, mat.step, 0, mat.cols * mat.elemSize(), 0, 0, NULL, NULL);
226 clFlush((cl_command_queue)mat.clCxt->oclCommandQueue());
230 devData = (cl_mem)mat.data;
233 clEnqueueCopyBufferToImage((cl_command_queue)mat.clCxt->oclCommandQueue(), devData, texture, 0, origin, region, 0, NULL, 0);
234 if ((mat.cols * mat.elemSize() != mat.step))
236 clFlush((cl_command_queue)mat.clCxt->oclCommandQueue());
237 clReleaseMemObject(devData);
244 GCC_DIAG_ON(deprecated-declarations)
247 Ptr<TextureCL> bindTexturePtr(const oclMat &mat)
249 return Ptr<TextureCL>(new TextureCL(bindTexture(mat), mat.rows, mat.cols, mat.type()));
251 void releaseTexture(cl_mem& texture)
256 bool support_image2d(Context *clCxt)
258 static const char * _kernel_string = "__kernel void test_func(image2d_t img) {}";
259 static bool _isTested = false;
260 static bool _support = false;
267 cv::ocl::openCLGetKernelFromSource(clCxt, &_kernel_string, "test_func");
271 catch (const cv::Exception& e)
279 // throw e once again