2121#include " qgsfeedback.h"
2222#include " qgsogrutils.h"
2323#include < QFile>
24+ #include < QDebug>
2425#include < QFileInfo>
26+ #include < iterator>
2527
2628#ifdef HAVE_OPENCL
27- #ifdef __APPLE__
28- #include < OpenCL/opencl.h>
29- #else
29+ #include < CL/cl.hpp>
3030#include < CL/cl.h>
3131#endif
32- #endif
3332
3433
3534QgsNineCellFilter::QgsNineCellFilter ( const QString &inputFile, const QString &outputFile, const QString &outputFormat )
@@ -97,45 +96,75 @@ int QgsNineCellFilter::processRaster( QgsFeedback *feedback )
9796 float *resultLine = ( float * ) CPLMalloc ( sizeof ( float ) * xSize );
9897
9998#ifdef HAVE_OPENCL
100- // TODO: move to utils and check for errors
99+
100+ cl_int errorCode = 0 ;
101101
102102 // Get platform and device information
103- cl_platform_id platform_id = NULL ;
104- cl_device_id device_id = NULL ;
105- cl_uint ret_num_devices;
106- cl_uint ret_num_platforms;
107- cl_int ret = clGetPlatformIDs ( 1 , &platform_id, &ret_num_platforms );
108- ret = clGetDeviceIDs ( platform_id, CL_DEVICE_TYPE_ALL, 1 ,
109- &device_id, &ret_num_devices );
110-
111- // Create an OpenCL context
112- cl_context context = clCreateContext ( NULL , 1 , &device_id, NULL , NULL , &ret );
113-
114- // Create a command queue
115- cl_command_queue command_queue = clCreateCommandQueue ( context, device_id, 0 , &ret );
116-
117- // Create memory buffers on the device for each vector
118- cl_mem scanLine1_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
119- sizeof ( float ) * ( xSize + 2 ), NULL , &ret );
120- cl_mem scanLine2_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
121- sizeof ( float ) * ( xSize + 2 ), NULL , &ret );
122- cl_mem scanLine3_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
123- sizeof ( float ) * ( xSize + 2 ), NULL , &ret );
103+ // cl_platform_id platform_id = NULL;
104+ // cl_device_id device_id = NULL;
105+ // cl_uint ret_num_devices;
106+ // cl_uint ret_num_platforms;
107+ // cl_int ret = clGetPlatformIDs( 1, &platform_id, &ret_num_platforms );
108+ // ret = clGetDeviceIDs( platform_id, CL_DEVICE_TYPE_ALL, 1,
109+ // &device_id, &ret_num_devices );
110+
111+ // // Create an OpenCL context
112+ // cl_context context = clCreateContext( NULL, 1, &device_id, NULL, NULL, &ret );
113+
114+ // // Create a command queue
115+ // cl_command_queue command_queue = clCreateCommandQueue( context, device_id, 0, &ret );
116+
117+ // // Create memory buffers on the device for each vector
118+ // cl_mem scanLine1Buffer = clCreateBuffer( context, CL_MEM_READ_ONLY,
119+ // sizeof( float ) * ( xSize + 2 ), NULL, &ret );
120+ // cl_mem scanLine2Buffer = clCreateBuffer( context, CL_MEM_READ_ONLY,
121+ // sizeof( float ) * ( xSize + 2 ), NULL, &ret );
122+ // cl_mem scanLine3Buffer = clCreateBuffer( context, CL_MEM_READ_ONLY,
123+ // sizeof( float ) * ( xSize + 2 ), NULL, &ret );
124124
125125 // TODO: constants
126- cl_mem inputNodataValue_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
127- sizeof ( float ), NULL , &ret );
128- cl_mem outputNodataValue_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
129- sizeof ( float ), NULL , &ret );
130- cl_mem zFactor_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
131- sizeof ( double ), NULL , &ret );
132- cl_mem cellSizeX_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
133- sizeof ( double ), NULL , &ret );
134- cl_mem cellSizeY_mem_obj = clCreateBuffer ( context, CL_MEM_READ_ONLY,
135- sizeof ( double ), NULL , &ret );
136126
137- cl_mem resultLine_mem_obj = clCreateBuffer ( context, CL_MEM_WRITE_ONLY,
138- sizeof ( float ) * xSize, NULL , &ret );
127+
128+ // cl_mem inputNodataValueBuffer = clCreateBuffer( context, CL_MEM_READ_ONLY,
129+ // sizeof( float ), NULL, &ret );
130+ // cl_mem outputNodataValueBuffer = clCreateBuffer( context, CL_MEM_READ_ONLY,
131+ // sizeof( float ), NULL, &ret );
132+ // cl_mem zFactorBuffer = clCreateBuffer( context, CL_MEM_READ_ONLY,
133+ // sizeof( double ), NULL, &ret );
134+ // cl_mem cellSizeXBuffer = clCreateBuffer( context, CL_MEM_READ_ONLY,
135+ // sizeof( double ), NULL, &ret );
136+ // cl_mem cellSizeYBuffer = clCreateBuffer( context, CL_MEM_READ_ONLY,
137+ // sizeof( double ), NULL, &ret );
138+
139+ // cl_mem resultLineBuffer = clCreateBuffer( context, CL_MEM_WRITE_ONLY,
140+ // sizeof( float ) * xSize, NULL, &ret );
141+
142+ std::vector<double > rasterParams;
143+
144+ rasterParams.push_back ( mInputNodataValue );
145+ rasterParams.push_back ( mOutputNodataValue );
146+ rasterParams.push_back ( mZFactor );
147+ rasterParams.push_back ( mCellSizeX );
148+ rasterParams.push_back ( mCellSizeY );
149+
150+ // cl::Buffer inputNodataValueBuffer( CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
151+ // sizeof( float ), mInputNodataValue , &ret );
152+ // cl::Buffer outputNodataValueBuffer( CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
153+ // sizeof( float ), mOutputNodataValue, &ret );
154+ // cl::Buffer zFactorBuffer ( CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
155+ // sizeof( double ), mZFactor, &ret );
156+ // cl::Buffer cellSizeXBuffer ( CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
157+ // sizeof( double ), mCellSizeX, &ret );
158+ // cl::Buffer cellSizeYBuffer ( CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR,
159+ // sizeof( double ), mCellSizeY, &ret );
160+
161+ cl::Buffer rasterParamsBuffer ( std::begin ( rasterParams ), std::end ( rasterParams ), true , false , &errorCode );
162+
163+ cl::Buffer resultLineBuffer ( CL_MEM_WRITE_ONLY, sizeof ( float ) * xSize, nullptr , &errorCode );
164+
165+ cl::Buffer scanLine1Buffer ( CL_MEM_READ_ONLY, sizeof ( float ) * ( xSize + 2 ), nullptr , &errorCode );
166+ cl::Buffer scanLine2Buffer ( CL_MEM_READ_ONLY, sizeof ( float ) * ( xSize + 2 ), nullptr , &errorCode );
167+ cl::Buffer scanLine3Buffer ( CL_MEM_READ_ONLY, sizeof ( float ) * ( xSize + 2 ), nullptr , &errorCode );
139168
140169
141170 char *source_str = new char [QFileInfo ( " /home/ale/dev/QGIS/src/analysis/raster/slope.cl" ).size () + 1 ];
@@ -148,36 +177,23 @@ int QgsNineCellFilter::processRaster( QgsFeedback *feedback )
148177 file.close ();
149178
150179 // Create a program from the kernel source
180+ cl::Program program ( source_str, true , &errorCode );
151181
152- Q_ASSERT ( ret == 0 );
153-
154- size_t source_size = strlen ( source_str );
155- cl_program program = clCreateProgramWithSource ( context, 1 ,
156- ( const char ** )&source_str, ( const size_t * )&source_size, &ret );
157-
158- // Build the program
159- ret = clBuildProgram ( program, 1 , &device_id, NULL , NULL , NULL );
160-
161- if ( ret != 0 )
182+ auto buildInfo = program.getBuildInfo <CL_PROGRAM_BUILD_LOG>( cl::Device::getDefault (), &errorCode );
183+ for ( auto &pair : buildInfo )
162184 {
163- char *program_log;
164- size_t log_size;
165- /* Find size of log and print to std output */
166- clGetProgramBuildInfo ( program, device_id, CL_PROGRAM_BUILD_LOG,
167- 0 , NULL , &log_size );
168- program_log = ( char * ) malloc ( log_size + 1 );
169- program_log[log_size] = ' \0 ' ;
170- clGetProgramBuildInfo ( program, device_id, CL_PROGRAM_BUILD_LOG,
171- log_size + 1 , program_log, NULL );
172- QgsDebugMsgLevel ( QStringLiteral ( " Error building OpenCL program: %1" ).arg ( program_log ), 1 );
173- free ( program_log );
185+ qDebug () << pair;
174186 }
175187
176-
177188 // Create the OpenCL kernel
178- cl_kernel kernel = clCreateKernel ( program, " processNineCellWindow" , &ret );
179-
180- Q_ASSERT ( ret == 0 );
189+ auto kernel =
190+ cl::make_kernel <
191+ cl::Buffer &,
192+ cl::Buffer &,
193+ cl::Buffer &,
194+ cl::Buffer &,
195+ cl::Buffer &
196+ > ( program, " processNineCellWindow" );
181197
182198#endif
183199
@@ -239,51 +255,69 @@ int QgsNineCellFilter::processRaster( QgsFeedback *feedback )
239255
240256#ifdef HAVE_OPENCL
241257 // Copy the scan lines to their respective memory buffers
242- ret = clEnqueueWriteBuffer ( command_queue, scanLine1_mem_obj, CL_TRUE, 0 ,
243- sizeof ( float ) * ( xSize + 2 ), scanLine1, 0 , NULL , NULL );
244- ret = clEnqueueWriteBuffer ( command_queue, scanLine2_mem_obj, CL_TRUE, 0 ,
245- sizeof ( float ) * ( xSize + 2 ), scanLine2, 0 , NULL , NULL );
246- ret = clEnqueueWriteBuffer ( command_queue, scanLine3_mem_obj, CL_TRUE, 0 ,
247- sizeof ( float ) * ( xSize + 2 ), scanLine3, 0 , NULL , NULL );
248-
249- ret = clEnqueueWriteBuffer ( command_queue, inputNodataValue_mem_obj, CL_TRUE, 0 ,
250- sizeof ( float ), &mInputNodataValue , 0 , NULL , NULL );
251- ret = clEnqueueWriteBuffer ( command_queue, outputNodataValue_mem_obj, CL_TRUE, 0 ,
252- sizeof ( float ), &mOutputNodataValue , 0 , NULL , NULL );
253- ret = clEnqueueWriteBuffer ( command_queue, zFactor_mem_obj, CL_TRUE, 0 ,
254- sizeof ( double ), &mZFactor , 0 , NULL , NULL );
255- ret = clEnqueueWriteBuffer ( command_queue, cellSizeX_mem_obj, CL_TRUE, 0 ,
256- sizeof ( double ), &mCellSizeX , 0 , NULL , NULL );
257- ret = clEnqueueWriteBuffer ( command_queue, cellSizeY_mem_obj, CL_TRUE, 0 ,
258- sizeof ( double ), &mCellSizeY , 0 , NULL , NULL );
259-
260-
261- // Set the arguments of the kernel
262- ret = ret || clSetKernelArg ( kernel, 0 , sizeof ( cl_mem ), ( void * )&scanLine1_mem_obj );
263- ret = ret || clSetKernelArg ( kernel, 1 , sizeof ( cl_mem ), ( void * )&scanLine2_mem_obj );
264- ret = ret || clSetKernelArg ( kernel, 2 , sizeof ( cl_mem ), ( void * )&scanLine3_mem_obj );
265- ret = ret || clSetKernelArg ( kernel, 3 , sizeof ( cl_mem ), ( void * )&resultLine_mem_obj );
266- ret = ret || clSetKernelArg ( kernel, 4 , sizeof ( cl_mem ), ( void * )&inputNodataValue_mem_obj );
267- ret = ret || clSetKernelArg ( kernel, 5 , sizeof ( cl_mem ), ( void * )&outputNodataValue_mem_obj );
268- ret = ret || clSetKernelArg ( kernel, 6 , sizeof ( cl_mem ), ( void * )&zFactor_mem_obj );
269- ret = ret || clSetKernelArg ( kernel, 7 , sizeof ( cl_mem ), ( void * )&cellSizeX_mem_obj );
270- ret = ret || clSetKernelArg ( kernel, 8 , sizeof ( cl_mem ), ( void * )&cellSizeY_mem_obj );
271-
272- Q_ASSERT ( ret == 0 );
273-
274- // Execute the OpenCL kernel on the scan line
275- size_t global_item_size = xSize; // Process the entire lists
276- // size_t local_item_size = 64; // Process in groups of 64 (or NULL for auto)
277- // ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
278- // &global_item_size, &local_item_size, 0, NULL, NULL);
279- ret = clEnqueueNDRangeKernel ( command_queue, kernel, 1 , NULL ,
280- &global_item_size, NULL , 0 , NULL , NULL );
281-
282- Q_ASSERT ( ret == 0 );
283-
284- ret = clEnqueueReadBuffer ( command_queue, resultLine_mem_obj, CL_TRUE, 0 ,
285- xSize * sizeof ( float ), resultLine, 0 , NULL , NULL );
286-
258+ // ret = clEnqueueWriteBuffer( command_queue, scanLine1Buffer, CL_TRUE, 0,
259+ // sizeof( float ) * ( xSize + 2 ), scanLine1, 0, NULL, NULL );
260+ // ret = clEnqueueWriteBuffer( command_queue, scanLine2Buffer, CL_TRUE, 0,
261+ // sizeof( float ) * ( xSize + 2 ), scanLine2, 0, NULL, NULL );
262+ // ret = clEnqueueWriteBuffer( command_queue, scanLine3Buffer, CL_TRUE, 0,
263+ // sizeof( float ) * ( xSize + 2 ), scanLine3, 0, NULL, NULL );
264+
265+ // ret = clEnqueueWriteBuffer( command_queue, inputNodataValueBuffer, CL_TRUE, 0,
266+ // sizeof( float ), &mInputNodataValue, 0, NULL, NULL );
267+ // ret = clEnqueueWriteBuffer( command_queue, outputNodataValueBuffer, CL_TRUE, 0,
268+ // sizeof( float ), &mOutputNodataValue, 0, NULL, NULL );
269+ // ret = clEnqueueWriteBuffer( command_queue, zFactorBuffer, CL_TRUE, 0,
270+ // sizeof( double ), &mZFactor, 0, NULL, NULL );
271+ // ret = clEnqueueWriteBuffer( command_queue, cellSizeXBuffer, CL_TRUE, 0,
272+ // sizeof( double ), &mCellSizeX, 0, NULL, NULL );
273+ // ret = clEnqueueWriteBuffer( command_queue, cellSizeYBuffer, CL_TRUE, 0,
274+ // sizeof( double ), &mCellSizeY, 0, NULL, NULL );
275+
276+
277+ // // Set the arguments of the kernel
278+ // ret = ret || clSetKernelArg( kernel, 0, sizeof( cl_mem ), ( void * )&scanLine1Buffer );
279+ // ret = ret || clSetKernelArg( kernel, 1, sizeof( cl_mem ), ( void * )&scanLine2Buffer );
280+ // ret = ret || clSetKernelArg( kernel, 2, sizeof( cl_mem ), ( void * )&scanLine3Buffer );
281+ // ret = ret || clSetKernelArg( kernel, 3, sizeof( cl_mem ), ( void * )&resultLineBuffer );
282+ // ret = ret || clSetKernelArg( kernel, 4, sizeof( cl_mem ), ( void * )&inputNodataValueBuffer );
283+ // ret = ret || clSetKernelArg( kernel, 5, sizeof( cl_mem ), ( void * )&outputNodataValueBuffer );
284+ // ret = ret || clSetKernelArg( kernel, 6, sizeof( cl_mem ), ( void * )&zFactorBuffer );
285+ // ret = ret || clSetKernelArg( kernel, 7, sizeof( cl_mem ), ( void * )&cellSizeXBuffer );
286+ // ret = ret || clSetKernelArg( kernel, 8, sizeof( cl_mem ), ( void * )&cellSizeYBuffer );
287+
288+
289+ errorCode = cl::enqueueWriteBuffer ( scanLine1Buffer, CL_TRUE, 0 ,
290+ sizeof ( float ) * ( xSize + 2 ), scanLine1 );
291+ errorCode = cl::enqueueWriteBuffer ( scanLine2Buffer, CL_TRUE, 0 ,
292+ sizeof ( float ) * ( xSize + 2 ), scanLine2 );
293+ errorCode = cl::enqueueWriteBuffer ( scanLine3Buffer, CL_TRUE, 0 ,
294+ sizeof ( float ) * ( xSize + 2 ), scanLine3 );
295+
296+
297+ kernel ( cl::EnqueueArgs (
298+ cl::NDRange ( xSize )
299+ ),
300+ scanLine1Buffer,
301+ scanLine2Buffer,
302+ scanLine3Buffer,
303+ resultLineBuffer,
304+ rasterParamsBuffer
305+ );
306+ // // Execute the OpenCL kernel on the scan line
307+ // size_t global_item_size = xSize; // Process the entire lists
308+ // //size_t local_item_size = 64; // Process in groups of 64 (or NULL for auto)
309+ // //ret = clEnqueueNDRangeKernel(command_queue, kernel, 1, NULL,
310+ // // &global_item_size, &local_item_size, 0, NULL, NULL);
311+ // ret = clEnqueueNDRangeKernel( command_queue, kernel, 1, NULL,
312+ // &global_item_size, NULL, 0, NULL, NULL );
313+
314+ // Q_ASSERT( ret == 0 );
315+
316+ // const cl_command_queue command_queue = cl::CommandQueue::getDefault()();
317+ // ret = clEnqueueReadBuffer( command_queue , resultLineBuffer(), CL_TRUE, 0,
318+ // xSize * sizeof( float ), resultLine, 0, NULL, NULL );
319+
320+ cl::enqueueReadBuffer ( resultLineBuffer, CL_TRUE, 0 , xSize * sizeof ( float ), resultLine );
287321
288322 if ( GDALRasterIO ( outputRasterBand, GF_Write, 0 , i, xSize, 1 , resultLine, xSize, 1 , GDT_Float32, 0 , 0 ) != CE_None )
289323 {
@@ -292,23 +326,6 @@ int QgsNineCellFilter::processRaster( QgsFeedback *feedback )
292326
293327 }
294328
295- // Clean up
296- // ret = clFlush( command_queue );
297- // ret = clFinish( command_queue );
298- ret = clReleaseKernel ( kernel );
299- ret = clReleaseProgram ( program );
300- ret = clReleaseMemObject ( scanLine1_mem_obj );
301- ret = clReleaseMemObject ( scanLine2_mem_obj );
302- ret = clReleaseMemObject ( scanLine3_mem_obj );
303- ret = clReleaseMemObject ( resultLine_mem_obj );
304- ret = clReleaseMemObject ( inputNodataValue_mem_obj );
305- ret = clReleaseMemObject ( outputNodataValue_mem_obj );
306- ret = clReleaseMemObject ( zFactor_mem_obj );
307- ret = clReleaseMemObject ( cellSizeX_mem_obj );
308- ret = clReleaseMemObject ( cellSizeY_mem_obj );
309- ret = clReleaseCommandQueue ( command_queue );
310- ret = clReleaseContext ( context );
311-
312329#else
313330
314331 // j is the x axis index, skip 0 and last cell that hve been filled with nodata
0 commit comments