Skip to content

Commit a7ef072

Browse files
elpasoelpaso
committed
Try to avoid crash on intel haswell
1 parent b6de8f1 commit a7ef072

File tree

4 files changed

+214
-97
lines changed

4 files changed

+214
-97
lines changed

‎src/analysis/raster/qgsninecellfilter.cpp‎

Lines changed: 115 additions & 94 deletions
Original file line numberDiff line numberDiff line change
@@ -50,33 +50,43 @@ int QgsNineCellFilter::processRaster( QgsFeedback *feedback )
5050
.arg( openClProgramBaseName( ) ), QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Info );
5151
return processRasterGPU( source, feedback );
5252
}
53-
catch ( cl::BuildError e )
53+
catch ( cl::BuildError &e )
5454
{
5555
cl::BuildLogType build_logs = e.getBuildLog();
5656
QString build_log;
5757
if ( build_logs.size() > 0 )
5858
build_log = QString::fromStdString( build_logs[0].second );
5959
else
6060
build_log = QObject::tr( "Build logs not available!" );
61-
QgsMessageLog::logMessage( QObject::tr( "Error building OpenCL program: %1" )
62-
.arg( build_log ), QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );
61+
QString err = QObject::tr( "Error building OpenCL program: %1" )
62+
.arg( build_log );
63+
QgsMessageLog::logMessage( err, QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );
64+
throw QgsProcessingException( err );
6365
}
64-
catch ( cl::Error e )
66+
catch ( cl::Error &e )
6567
{
66-
QgsMessageLog::logMessage( QObject::tr( "Error %1 running OpenCL program in %2" )
67-
.arg( QString::number( e.err() ), QString::fromStdString( e.what() ) ), QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );
68-
68+
QString err = QObject::tr( "Error %1 running OpenCL program in %2" )
69+
.arg( QgsOpenClUtils::errorText( e.err() ), QString::fromStdString( e.what() ) );
70+
QgsMessageLog::logMessage( err, QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );
71+
throw QgsProcessingException( err );
6972
}
7073
}
7174
else
7275
{
73-
QgsMessageLog::logMessage( QObject::tr( "Error loading OpenCL program sources" ),
76+
QString err = QObject::tr( "Error loading OpenCL program sources" );
77+
QgsMessageLog::logMessage( err,
7478
QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );
75-
79+
throw QgsProcessingException( err );
7680
}
7781
}
78-
#endif
82+
else
83+
{
84+
return processRasterCPU( feedback );
85+
}
86+
return 1;
87+
#else
7988
return processRasterCPU( feedback );
89+
#endif
8090
}
8191

8292
gdal::dataset_unique_ptr QgsNineCellFilter::openInputFile( int &nCellsX, int &nCellsY )
@@ -230,106 +240,117 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
230240

231241
addExtraRasterParams( rasterParams );
232242

233-
cl::Buffer rasterParamsBuffer( rasterParams.begin(), rasterParams.end(), true, false, &errorCode );
234-
cl::Buffer scanLine1Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
235-
cl::Buffer scanLine2Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
236-
cl::Buffer scanLine3Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
237-
cl::Buffer resultLineBuffer( CL_MEM_WRITE_ONLY, sizeof( float ) * xSize, nullptr, &errorCode );
238-
239-
// Create a program from the kernel source
240-
cl::Program program( source.toStdString() );
241-
// Uuse CL 1.1 for compatibility with older libs
242-
program.build( "-cl-std=CL1.1" );
243-
244-
// Create the OpenCL kernel
245-
auto kernel = cl::KernelFunctor <
246-
cl::Buffer &,
247-
cl::Buffer &,
248-
cl::Buffer &,
249-
cl::Buffer &,
250-
cl::Buffer &
251-
> ( program, "processNineCellWindow" );
252-
253-
//values outside the layer extent (if the 3x3 window is on the border) are sent to the processing method as (input) nodata values
254-
for ( int i = 0; i < ySize; ++i )
243+
try
255244
{
256-
if ( feedback && feedback->isCanceled() )
257-
{
258-
break;
259-
}
260245

261-
if ( feedback )
246+
cl::Buffer rasterParamsBuffer( rasterParams.begin(), rasterParams.end(), true, false, &errorCode );
247+
cl::Buffer scanLine1Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
248+
cl::Buffer scanLine2Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
249+
cl::Buffer scanLine3Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
250+
cl::Buffer resultLineBuffer( CL_MEM_WRITE_ONLY, sizeof( float ) * xSize, nullptr, &errorCode );
251+
252+
// Create a program from the kernel source
253+
cl::Program program( source.toStdString() );
254+
// Use CL 1.1 for compatibility with older libs
255+
program.build( "-cl-std=CL1.1" );
256+
257+
// Create the OpenCL kernel
258+
auto kernel = cl::KernelFunctor <
259+
cl::Buffer &,
260+
cl::Buffer &,
261+
cl::Buffer &,
262+
cl::Buffer &,
263+
cl::Buffer &
264+
> ( program, "processNineCellWindow" );
265+
266+
//values outside the layer extent (if the 3x3 window is on the border) are sent to the processing method as (input) nodata values
267+
for ( int i = 0; i < ySize; ++i )
262268
{
263-
feedback->setProgress( 100.0 * static_cast< double >( i ) / ySize );
264-
}
269+
if ( feedback && feedback->isCanceled() )
270+
{
271+
break;
272+
}
265273

266-
if ( i == 0 )
267-
{
268-
//fill scanline 1 with (input) nodata for the values above the first row and feed scanline2 with the first row
269-
for ( int a = 0; a < xSize + 2 ; ++a )
274+
if ( feedback )
270275
{
271-
scanLine1[a] = mInputNodataValue;
276+
feedback->setProgress( 100.0 * static_cast< double >( i ) / ySize );
272277
}
273-
// Read scanline2
274-
if ( GDALRasterIO( rasterBand, GF_Read, 0, 0, xSize, 1, &scanLine2[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
278+
279+
if ( i == 0 )
275280
{
276-
QgsDebugMsg( "Raster IO Error" );
281+
//fill scanline 1 with (input) nodata for the values above the first row and feed scanline2 with the first row
282+
for ( int a = 0; a < xSize + 2 ; ++a )
283+
{
284+
scanLine1[a] = mInputNodataValue;
285+
}
286+
// Read scanline2
287+
if ( GDALRasterIO( rasterBand, GF_Read, 0, 0, xSize, 1, &scanLine2[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
288+
{
289+
QgsDebugMsg( "Raster IO Error" );
290+
}
291+
}
292+
else
293+
{
294+
//normally fetch only scanLine3 and release scanline 1 if we move forward one row
295+
CPLFree( scanLine1 );
296+
scanLine1 = scanLine2;
297+
scanLine2 = scanLine3;
298+
scanLine3 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) );
277299
}
278-
}
279-
else
280-
{
281-
//normally fetch only scanLine3 and release scanline 1 if we move forward one row
282-
CPLFree( scanLine1 );
283-
scanLine1 = scanLine2;
284-
scanLine2 = scanLine3;
285-
scanLine3 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) );
286-
}
287300

288-
// Read scanline 3
289-
if ( i == ySize - 1 ) //fill the row below the bottom with nodata values
290-
{
291-
for ( int a = 0; a < xSize + 2; ++a )
301+
// Read scanline 3
302+
if ( i == ySize - 1 ) //fill the row below the bottom with nodata values
292303
{
293-
scanLine3[a] = mInputNodataValue;
304+
for ( int a = 0; a < xSize + 2; ++a )
305+
{
306+
scanLine3[a] = mInputNodataValue;
307+
}
294308
}
295-
}
296-
else
297-
{
298-
if ( GDALRasterIO( rasterBand, GF_Read, 0, i + 1, xSize, 1, &scanLine3[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
309+
else
310+
{
311+
if ( GDALRasterIO( rasterBand, GF_Read, 0, i + 1, xSize, 1, &scanLine3[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
312+
{
313+
QgsDebugMsg( "Raster IO Error" );
314+
}
315+
}
316+
// Set first and last extra colums to nodata
317+
scanLine1[0] = scanLine1[xSize + 1] = mInputNodataValue;
318+
scanLine2[0] = scanLine2[xSize + 1] = mInputNodataValue;
319+
scanLine3[0] = scanLine3[xSize + 1] = mInputNodataValue;
320+
321+
errorCode = cl::enqueueWriteBuffer( scanLine1Buffer, CL_TRUE, 0,
322+
sizeof( float ) * ( xSize + 2 ), scanLine1 );
323+
errorCode = cl::enqueueWriteBuffer( scanLine2Buffer, CL_TRUE, 0,
324+
sizeof( float ) * ( xSize + 2 ), scanLine2 );
325+
errorCode = cl::enqueueWriteBuffer( scanLine3Buffer, CL_TRUE, 0,
326+
sizeof( float ) * ( xSize + 2 ), scanLine3 );
327+
328+
kernel( cl::EnqueueArgs(
329+
cl::NDRange( xSize )
330+
),
331+
scanLine1Buffer,
332+
scanLine2Buffer,
333+
scanLine3Buffer,
334+
resultLineBuffer,
335+
rasterParamsBuffer
336+
);
337+
338+
cl::enqueueReadBuffer( resultLineBuffer, CL_TRUE, 0, xSize * sizeof( float ), resultLine );
339+
340+
if ( GDALRasterIO( outputRasterBand, GF_Write, 0, i, xSize, 1, resultLine, xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
299341
{
300342
QgsDebugMsg( "Raster IO Error" );
301343
}
302-
}
303-
// Set first and last extra colums to nodata
304-
scanLine1[0] = scanLine1[xSize + 1] = mInputNodataValue;
305-
scanLine2[0] = scanLine2[xSize + 1] = mInputNodataValue;
306-
scanLine3[0] = scanLine3[xSize + 1] = mInputNodataValue;
307-
308-
errorCode = cl::enqueueWriteBuffer( scanLine1Buffer, CL_TRUE, 0,
309-
sizeof( float ) * ( xSize + 2 ), scanLine1 );
310-
errorCode = cl::enqueueWriteBuffer( scanLine2Buffer, CL_TRUE, 0,
311-
sizeof( float ) * ( xSize + 2 ), scanLine2 );
312-
errorCode = cl::enqueueWriteBuffer( scanLine3Buffer, CL_TRUE, 0,
313-
sizeof( float ) * ( xSize + 2 ), scanLine3 );
314-
315344

316-
kernel( cl::EnqueueArgs(
317-
cl::NDRange( xSize )
318-
),
319-
scanLine1Buffer,
320-
scanLine2Buffer,
321-
scanLine3Buffer,
322-
resultLineBuffer,
323-
rasterParamsBuffer
324-
);
325-
326-
cl::enqueueReadBuffer( resultLineBuffer, CL_TRUE, 0, xSize * sizeof( float ), resultLine );
327-
328-
if ( GDALRasterIO( outputRasterBand, GF_Write, 0, i, xSize, 1, resultLine, xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
329-
{
330-
QgsDebugMsg( "Raster IO Error" );
331345
}
332-
346+
}
347+
catch ( cl::Error &e )
348+
{
349+
CPLFree( resultLine );
350+
CPLFree( scanLine1 );
351+
CPLFree( scanLine2 );
352+
CPLFree( scanLine3 );
353+
throw e;
333354
}
334355

335356
CPLFree( resultLine );

‎src/analysis/raster/slope.cl‎

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -1,5 +1,3 @@
1-
#pragma OPENCL EXTENSION cl_khr_fp64 : enable
2-
31
float calcFirstDer( float x11, float x21, float x31, float x12, float x22, float x32, float x13, float x23, float x33,
42
float mInputNodataValue, float mOutputNodataValue, float mZFactor, float mCellSize )
53
{
@@ -79,7 +77,7 @@ __kernel void processNineCellWindow( __global float *scanLine1,
7977
) {
8078

8179
// Get the index of the current element
82-
int i = get_global_id(0);
80+
const int i = get_global_id(0);
8381

8482
// Do the operation
8583
//return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX))

‎src/core/qgsopenclutils.cpp‎

Lines changed: 97 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -122,3 +122,100 @@ QString QgsOpenClUtils::buildLog( cl::BuildError &e )
122122
build_log = QString::fromStdString( build_logs[0].second );
123123
return build_log;
124124
}
125+
126+
QString QgsOpenClUtils::errorText( const int errorCode )
127+
{
128+
switch ( errorCode )
129+
{
130+
case 0: return QStringLiteral( "CL_SUCCESS" );
131+
case -1: return QStringLiteral( "CL_DEVICE_NOT_FOUND" );
132+
case -2: return QStringLiteral( "CL_DEVICE_NOT_AVAILABLE" );
133+
case -3: return QStringLiteral( "CL_COMPILER_NOT_AVAILABLE" );
134+
case -4: return QStringLiteral( "CL_MEM_OBJECT_ALLOCATION_FAILURE" );
135+
case -5: return QStringLiteral( "CL_OUT_OF_RESOURCES" );
136+
case -6: return QStringLiteral( "CL_OUT_OF_HOST_MEMORY" );
137+
case -7: return QStringLiteral( "CL_PROFILING_INFO_NOT_AVAILABLE" );
138+
case -8: return QStringLiteral( "CL_MEM_COPY_OVERLAP" );
139+
case -9: return QStringLiteral( "CL_IMAGE_FORMAT_MISMATCH" );
140+
case -10: return QStringLiteral( "CL_IMAGE_FORMAT_NOT_SUPPORTED" );
141+
case -12: return QStringLiteral( "CL_MAP_FAILURE" );
142+
case -13: return QStringLiteral( "CL_MISALIGNED_SUB_BUFFER_OFFSET" );
143+
case -14: return QStringLiteral( "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST" );
144+
case -15: return QStringLiteral( "CL_COMPILE_PROGRAM_FAILURE" );
145+
case -16: return QStringLiteral( "CL_LINKER_NOT_AVAILABLE" );
146+
case -17: return QStringLiteral( "CL_LINK_PROGRAM_FAILURE" );
147+
case -18: return QStringLiteral( "CL_DEVICE_PARTITION_FAILED" );
148+
case -19: return QStringLiteral( "CL_KERNEL_ARG_INFO_NOT_AVAILABLE" );
149+
case -30: return QStringLiteral( "CL_INVALID_VALUE" );
150+
case -31: return QStringLiteral( "CL_INVALID_DEVICE_TYPE" );
151+
case -32: return QStringLiteral( "CL_INVALID_PLATFORM" );
152+
case -33: return QStringLiteral( "CL_INVALID_DEVICE" );
153+
case -34: return QStringLiteral( "CL_INVALID_CONTEXT" );
154+
case -35: return QStringLiteral( "CL_INVALID_QUEUE_PROPERTIES" );
155+
case -36: return QStringLiteral( "CL_INVALID_COMMAND_QUEUE" );
156+
case -37: return QStringLiteral( "CL_INVALID_HOST_PTR" );
157+
case -38: return QStringLiteral( "CL_INVALID_MEM_OBJECT" );
158+
case -39: return QStringLiteral( "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR" );
159+
case -40: return QStringLiteral( "CL_INVALID_IMAGE_SIZE" );
160+
case -41: return QStringLiteral( "CL_INVALID_SAMPLER" );
161+
case -42: return QStringLiteral( "CL_INVALID_BINARY" );
162+
case -43: return QStringLiteral( "CL_INVALID_BUILD_OPTIONS" );
163+
case -44: return QStringLiteral( "CL_INVALID_PROGRAM" );
164+
case -45: return QStringLiteral( "CL_INVALID_PROGRAM_EXECUTABLE" );
165+
case -46: return QStringLiteral( "CL_INVALID_KERNEL_NAME" );
166+
case -47: return QStringLiteral( "CL_INVALID_KERNEL_DEFINITION" );
167+
case -48: return QStringLiteral( "CL_INVALID_KERNEL" );
168+
case -49: return QStringLiteral( "CL_INVALID_ARG_INDEX" );
169+
case -50: return QStringLiteral( "CL_INVALID_ARG_VALUE" );
170+
case -51: return QStringLiteral( "CL_INVALID_ARG_SIZE" );
171+
case -52: return QStringLiteral( "CL_INVALID_KERNEL_ARGS" );
172+
case -53: return QStringLiteral( "CL_INVALID_WORK_DIMENSION" );
173+
case -54: return QStringLiteral( "CL_INVALID_WORK_GROUP_SIZE" );
174+
case -55: return QStringLiteral( "CL_INVALID_WORK_ITEM_SIZE" );
175+
case -56: return QStringLiteral( "CL_INVALID_GLOBAL_OFFSET" );
176+
case -57: return QStringLiteral( "CL_INVALID_EVENT_WAIT_LIST" );
177+
case -58: return QStringLiteral( "CL_INVALID_EVENT" );
178+
case -59: return QStringLiteral( "CL_INVALID_OPERATION" );
179+
case -60: return QStringLiteral( "CL_INVALID_GL_OBJECT" );
180+
case -61: return QStringLiteral( "CL_INVALID_BUFFER_SIZE" );
181+
case -62: return QStringLiteral( "CL_INVALID_MIP_LEVEL" );
182+
case -63: return QStringLiteral( "CL_INVALID_GLOBAL_WORK_SIZE" );
183+
case -64: return QStringLiteral( "CL_INVALID_PROPERTY" );
184+
case -65: return QStringLiteral( "CL_INVALID_IMAGE_DESCRIPTOR" );
185+
case -66: return QStringLiteral( "CL_INVALID_COMPILER_OPTIONS" );
186+
case -67: return QStringLiteral( "CL_INVALID_LINKER_OPTIONS" );
187+
case -68: return QStringLiteral( "CL_INVALID_DEVICE_PARTITION_COUNT" );
188+
case -69: return QStringLiteral( "CL_INVALID_PIPE_SIZE" );
189+
case -70: return QStringLiteral( "CL_INVALID_DEVICE_QUEUE" );
190+
case -71: return QStringLiteral( "CL_INVALID_SPEC_ID" );
191+
case -72: return QStringLiteral( "CL_MAX_SIZE_RESTRICTION_EXCEEDED" );
192+
case -1002: return QStringLiteral( "CL_INVALID_D3D10_DEVICE_KHR" );
193+
case -1003: return QStringLiteral( "CL_INVALID_D3D10_RESOURCE_KHR" );
194+
case -1004: return QStringLiteral( "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR" );
195+
case -1005: return QStringLiteral( "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR" );
196+
case -1006: return QStringLiteral( "CL_INVALID_D3D11_DEVICE_KHR" );
197+
case -1007: return QStringLiteral( "CL_INVALID_D3D11_RESOURCE_KHR" );
198+
case -1008: return QStringLiteral( "CL_D3D11_RESOURCE_ALREADY_ACQUIRED_KHR" );
199+
case -1009: return QStringLiteral( "CL_D3D11_RESOURCE_NOT_ACQUIRED_KHR" );
200+
case -1010: return QStringLiteral( "CL_INVALID_DX9_MEDIA_ADAPTER_KHR" );
201+
case -1011: return QStringLiteral( "CL_INVALID_DX9_MEDIA_SURFACE_KHR" );
202+
case -1012: return QStringLiteral( "CL_DX9_MEDIA_SURFACE_ALREADY_ACQUIRED_KHR" );
203+
case -1013: return QStringLiteral( "CL_DX9_MEDIA_SURFACE_NOT_ACQUIRED_KHR" );
204+
case -1093: return QStringLiteral( "CL_INVALID_EGL_OBJECT_KHR" );
205+
case -1092: return QStringLiteral( "CL_EGL_RESOURCE_NOT_ACQUIRED_KHR" );
206+
case -1001: return QStringLiteral( "CL_PLATFORM_NOT_FOUND_KHR" );
207+
case -1057: return QStringLiteral( "CL_DEVICE_PARTITION_FAILED_EXT" );
208+
case -1058: return QStringLiteral( "CL_INVALID_PARTITION_COUNT_EXT" );
209+
case -1059: return QStringLiteral( "CL_INVALID_PARTITION_NAME_EXT" );
210+
case -1094: return QStringLiteral( "CL_INVALID_ACCELERATOR_INTEL" );
211+
case -1095: return QStringLiteral( "CL_INVALID_ACCELERATOR_TYPE_INTEL" );
212+
case -1096: return QStringLiteral( "CL_INVALID_ACCELERATOR_DESCRIPTOR_INTEL" );
213+
case -1097: return QStringLiteral( "CL_ACCELERATOR_TYPE_NOT_SUPPORTED_INTEL" );
214+
case -1000: return QStringLiteral( "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR" );
215+
case -1098: return QStringLiteral( "CL_INVALID_VA_API_MEDIA_ADAPTER_INTEL" );
216+
case -1099: return QStringLiteral( "CL_INVALID_VA_API_MEDIA_SURFACE_INTEL" );
217+
case -1100: return QStringLiteral( "CL_VA_API_MEDIA_SURFACE_ALREADY_ACQUIRED_INTEL" );
218+
case -1101: return QStringLiteral( "CL_VA_API_MEDIA_SURFACE_NOT_ACQUIRED_INTEL" );
219+
default: return QStringLiteral( "CL_UNKNOWN_ERROR" );
220+
}
221+
}

‎src/core/qgsopenclutils.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -44,6 +44,7 @@ class CORE_EXPORT QgsOpenClUtils
4444
static QString buildLog( cl::BuildError &e );
4545
static QString sourceFromPath( const QString &path );
4646
static QLatin1String LOGMESSAGE_TAG;
47+
static QString errorText( const int errorCode );
4748

4849
private:
4950
QgsOpenClUtils();

0 commit comments

Comments
 (0)