Skip to content

Commit a9f11fb

Browse files
elpasoelpaso
committedAug 8, 2018
Working! With C++ API
·
ltr-3_40final-3_4_0
1 parent 0c3cb68 commit a9f11fb

File tree

2 files changed

+153
-138
lines changed

2 files changed

+153
-138
lines changed
 

‎src/analysis/raster/qgsninecellfilter.cpp

Lines changed: 142 additions & 125 deletions
Original file line numberDiff line numberDiff line change
@@ -21,15 +21,14 @@
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

3534
QgsNineCellFilter::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

‎src/analysis/raster/slope.cl

Lines changed: 11 additions & 13 deletions
Original file line numberDiff line numberDiff line change
@@ -1,7 +1,7 @@
11
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
22

33
float calcFirstDer( float x11, float x21, float x31, float x12, float x22, float x32, float x13, float x23, float x33,
4-
float mInputNodataValue, float mOutputNodataValue, double mZFactor, double mCellSize )
4+
double mInputNodataValue, double mOutputNodataValue, double mZFactor, double mCellSize )
55
{
66
//the basic formula would be simple, but we need to test for nodata values...
77
//X: return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX));
@@ -72,14 +72,10 @@ float calcFirstDer( float x11, float x21, float x31, float x12, float x22, float
7272

7373

7474
__kernel void processNineCellWindow( __global float *scanLine1,
75-
__global float *scanLine2,
76-
__global float *scanLine3,
77-
__global float *resultLine,
78-
__global float *mInputNodataValue,
79-
__global float *mOutputNodataValue,
80-
__global double *mZFactor,
81-
__global double *mCellSizeX,
82-
__global double *mCellSizeY
75+
__global float *scanLine2,
76+
__global float *scanLine3,
77+
__global float *resultLine,
78+
__global double *rasterParams
8379
) {
8480

8581
// Get the index of the current element
@@ -90,17 +86,19 @@ __kernel void processNineCellWindow( __global float *scanLine1,
9086
float derX = calcFirstDer( scanLine1[i], scanLine2[i], scanLine3[i],
9187
scanLine1[i+1], scanLine2[i+1], scanLine3[i+1],
9288
scanLine1[i+2], scanLine2[i+2], scanLine3[i+2],
93-
*mInputNodataValue, *mOutputNodataValue, *mZFactor, *mCellSizeX
89+
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[3]
9490
);
9591
//return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * mCellSizeY));
9692
float derY = calcFirstDer( scanLine1[i+2], scanLine1[i+1], scanLine1[i],
9793
scanLine2[i+2], scanLine2[i+1], scanLine2[i],
9894
scanLine3[i+2], scanLine3[i+1], scanLine3[i],
99-
*mInputNodataValue, *mOutputNodataValue, *mZFactor, *mCellSizeY
95+
rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[4]
10096
);
101-
if ( derX == *mOutputNodataValue || derY == *mOutputNodataValue )
97+
98+
99+
if ( derX == rasterParams[1] || derY == rasterParams[1] )
102100
{
103-
resultLine[i] = *mOutputNodataValue;
101+
resultLine[i] = rasterParams[1];
104102
}
105103
else
106104
{

0 commit comments

Comments
 (0)
Please sign in to comment.