Skip to content

Commit

Permalink
Try to avoid crash on intel haswell
Browse files Browse the repository at this point in the history
  • Loading branch information
@elpaso
elpaso committed Aug 8, 2018
1 parent b6de8f1 commit a7ef072
Show file tree
Hide file tree
Showing 4 changed files with 214 additions and 97 deletions.
209 changes: 115 additions & 94 deletions src/analysis/raster/qgsninecellfilter.cpp
View file
Expand Up @@ -50,33 +50,43 @@ int QgsNineCellFilter::processRaster( QgsFeedback *feedback )
.arg( openClProgramBaseName( ) ), QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Info );
return processRasterGPU( source, feedback );
}
catch ( cl::BuildError e )
catch ( cl::BuildError &e )
{
cl::BuildLogType build_logs = e.getBuildLog();
QString build_log;
if ( build_logs.size() > 0 )
build_log = QString::fromStdString( build_logs[0].second );
else
build_log = QObject::tr( "Build logs not available!" );
QgsMessageLog::logMessage( QObject::tr( "Error building OpenCL program: %1" )
.arg( build_log ), QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );
QString err = QObject::tr( "Error building OpenCL program: %1" )
.arg( build_log );
QgsMessageLog::logMessage( err, QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );
throw QgsProcessingException( err );
}
catch ( cl::Error e )
catch ( cl::Error &e )
{
QgsMessageLog::logMessage( QObject::tr( "Error %1 running OpenCL program in %2" )
.arg( QString::number( e.err() ), QString::fromStdString( e.what() ) ), QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );

QString err = QObject::tr( "Error %1 running OpenCL program in %2" )
.arg( QgsOpenClUtils::errorText( e.err() ), QString::fromStdString( e.what() ) );
QgsMessageLog::logMessage( err, QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );
throw QgsProcessingException( err );
}
}
else
{
QgsMessageLog::logMessage( QObject::tr( "Error loading OpenCL program sources" ),
QString err = QObject::tr( "Error loading OpenCL program sources" );
QgsMessageLog::logMessage( err,
QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );

throw QgsProcessingException( err );
}
}
#endif
else
{
return processRasterCPU( feedback );
}
return 1;
#else
return processRasterCPU( feedback );
#endif
}

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

addExtraRasterParams( rasterParams );

cl::Buffer rasterParamsBuffer( rasterParams.begin(), rasterParams.end(), true, false, &errorCode );
cl::Buffer scanLine1Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
cl::Buffer scanLine2Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
cl::Buffer scanLine3Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
cl::Buffer resultLineBuffer( CL_MEM_WRITE_ONLY, sizeof( float ) * xSize, nullptr, &errorCode );

// Create a program from the kernel source
cl::Program program( source.toStdString() );
// Uuse CL 1.1 for compatibility with older libs
program.build( "-cl-std=CL1.1" );

// Create the OpenCL kernel
auto kernel = cl::KernelFunctor <
cl::Buffer &,
cl::Buffer &,
cl::Buffer &,
cl::Buffer &,
cl::Buffer &
> ( program, "processNineCellWindow" );

//values outside the layer extent (if the 3x3 window is on the border) are sent to the processing method as (input) nodata values
for ( int i = 0; i < ySize; ++i )
try
{
if ( feedback && feedback->isCanceled() )
{
break;
}

if ( feedback )
cl::Buffer rasterParamsBuffer( rasterParams.begin(), rasterParams.end(), true, false, &errorCode );
cl::Buffer scanLine1Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
cl::Buffer scanLine2Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
cl::Buffer scanLine3Buffer( CL_MEM_READ_ONLY, sizeof( float ) * ( xSize + 2 ), nullptr, &errorCode );
cl::Buffer resultLineBuffer( CL_MEM_WRITE_ONLY, sizeof( float ) * xSize, nullptr, &errorCode );

// Create a program from the kernel source
cl::Program program( source.toStdString() );
// Use CL 1.1 for compatibility with older libs
program.build( "-cl-std=CL1.1" );

// Create the OpenCL kernel
auto kernel = cl::KernelFunctor <
cl::Buffer &,
cl::Buffer &,
cl::Buffer &,
cl::Buffer &,
cl::Buffer &
> ( program, "processNineCellWindow" );

//values outside the layer extent (if the 3x3 window is on the border) are sent to the processing method as (input) nodata values
for ( int i = 0; i < ySize; ++i )
{
feedback->setProgress( 100.0 * static_cast< double >( i ) / ySize );
}
if ( feedback && feedback->isCanceled() )
{
break;
}

if ( i == 0 )
{
//fill scanline 1 with (input) nodata for the values above the first row and feed scanline2 with the first row
for ( int a = 0; a < xSize + 2 ; ++a )
if ( feedback )
{
scanLine1[a] = mInputNodataValue;
feedback->setProgress( 100.0 * static_cast< double >( i ) / ySize );
}
// Read scanline2
if ( GDALRasterIO( rasterBand, GF_Read, 0, 0, xSize, 1, &scanLine2[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )

if ( i == 0 )
{
QgsDebugMsg( "Raster IO Error" );
//fill scanline 1 with (input) nodata for the values above the first row and feed scanline2 with the first row
for ( int a = 0; a < xSize + 2 ; ++a )
{
scanLine1[a] = mInputNodataValue;
}
// Read scanline2
if ( GDALRasterIO( rasterBand, GF_Read, 0, 0, xSize, 1, &scanLine2[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
{
QgsDebugMsg( "Raster IO Error" );
}
}
else
{
//normally fetch only scanLine3 and release scanline 1 if we move forward one row
CPLFree( scanLine1 );
scanLine1 = scanLine2;
scanLine2 = scanLine3;
scanLine3 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) );
}
}
else
{
//normally fetch only scanLine3 and release scanline 1 if we move forward one row
CPLFree( scanLine1 );
scanLine1 = scanLine2;
scanLine2 = scanLine3;
scanLine3 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) );
}

// Read scanline 3
if ( i == ySize - 1 ) //fill the row below the bottom with nodata values
{
for ( int a = 0; a < xSize + 2; ++a )
// Read scanline 3
if ( i == ySize - 1 ) //fill the row below the bottom with nodata values
{
scanLine3[a] = mInputNodataValue;
for ( int a = 0; a < xSize + 2; ++a )
{
scanLine3[a] = mInputNodataValue;
}
}
}
else
{
if ( GDALRasterIO( rasterBand, GF_Read, 0, i + 1, xSize, 1, &scanLine3[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
else
{
if ( GDALRasterIO( rasterBand, GF_Read, 0, i + 1, xSize, 1, &scanLine3[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
{
QgsDebugMsg( "Raster IO Error" );
}
}
// Set first and last extra colums to nodata
scanLine1[0] = scanLine1[xSize + 1] = mInputNodataValue;
scanLine2[0] = scanLine2[xSize + 1] = mInputNodataValue;
scanLine3[0] = scanLine3[xSize + 1] = mInputNodataValue;

errorCode = cl::enqueueWriteBuffer( scanLine1Buffer, CL_TRUE, 0,
sizeof( float ) * ( xSize + 2 ), scanLine1 );
errorCode = cl::enqueueWriteBuffer( scanLine2Buffer, CL_TRUE, 0,
sizeof( float ) * ( xSize + 2 ), scanLine2 );
errorCode = cl::enqueueWriteBuffer( scanLine3Buffer, CL_TRUE, 0,
sizeof( float ) * ( xSize + 2 ), scanLine3 );

kernel( cl::EnqueueArgs(
cl::NDRange( xSize )
),
scanLine1Buffer,
scanLine2Buffer,
scanLine3Buffer,
resultLineBuffer,
rasterParamsBuffer
);

cl::enqueueReadBuffer( resultLineBuffer, CL_TRUE, 0, xSize * sizeof( float ), resultLine );

if ( GDALRasterIO( outputRasterBand, GF_Write, 0, i, xSize, 1, resultLine, xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
{
QgsDebugMsg( "Raster IO Error" );
}
}
// Set first and last extra colums to nodata
scanLine1[0] = scanLine1[xSize + 1] = mInputNodataValue;
scanLine2[0] = scanLine2[xSize + 1] = mInputNodataValue;
scanLine3[0] = scanLine3[xSize + 1] = mInputNodataValue;

errorCode = cl::enqueueWriteBuffer( scanLine1Buffer, CL_TRUE, 0,
sizeof( float ) * ( xSize + 2 ), scanLine1 );
errorCode = cl::enqueueWriteBuffer( scanLine2Buffer, CL_TRUE, 0,
sizeof( float ) * ( xSize + 2 ), scanLine2 );
errorCode = cl::enqueueWriteBuffer( scanLine3Buffer, CL_TRUE, 0,
sizeof( float ) * ( xSize + 2 ), scanLine3 );


kernel( cl::EnqueueArgs(
cl::NDRange( xSize )
),
scanLine1Buffer,
scanLine2Buffer,
scanLine3Buffer,
resultLineBuffer,
rasterParamsBuffer
);

cl::enqueueReadBuffer( resultLineBuffer, CL_TRUE, 0, xSize * sizeof( float ), resultLine );

if ( GDALRasterIO( outputRasterBand, GF_Write, 0, i, xSize, 1, resultLine, xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
{
QgsDebugMsg( "Raster IO Error" );
}

}
catch ( cl::Error &e )
{
CPLFree( resultLine );
CPLFree( scanLine1 );
CPLFree( scanLine2 );
CPLFree( scanLine3 );
throw e;
}

CPLFree( resultLine );
Expand Down
4 changes: 1 addition & 3 deletions src/analysis/raster/slope.cl
View file
@@ -1,5 +1,3 @@
#pragma OPENCL EXTENSION cl_khr_fp64 : enable

float calcFirstDer( float x11, float x21, float x31, float x12, float x22, float x32, float x13, float x23, float x33,
float mInputNodataValue, float mOutputNodataValue, float mZFactor, float mCellSize )
{
Expand Down Expand Up @@ -79,7 +77,7 @@ __kernel void processNineCellWindow( __global float *scanLine1,
) {

// Get the index of the current element
int i = get_global_id(0);
const int i = get_global_id(0);

// Do the operation
//return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX))
Expand Down
97 changes: 97 additions & 0 deletions src/core/qgsopenclutils.cpp
View file
Expand Up @@ -122,3 +122,100 @@ QString QgsOpenClUtils::buildLog( cl::BuildError &e )
build_log = QString::fromStdString( build_logs[0].second );
return build_log;
}

QString QgsOpenClUtils::errorText( const int errorCode )
{
switch ( errorCode )
{
case 0: return QStringLiteral( "CL_SUCCESS" );
case -1: return QStringLiteral( "CL_DEVICE_NOT_FOUND" );
case -2: return QStringLiteral( "CL_DEVICE_NOT_AVAILABLE" );
case -3: return QStringLiteral( "CL_COMPILER_NOT_AVAILABLE" );
case -4: return QStringLiteral( "CL_MEM_OBJECT_ALLOCATION_FAILURE" );
case -5: return QStringLiteral( "CL_OUT_OF_RESOURCES" );
case -6: return QStringLiteral( "CL_OUT_OF_HOST_MEMORY" );
case -7: return QStringLiteral( "CL_PROFILING_INFO_NOT_AVAILABLE" );
case -8: return QStringLiteral( "CL_MEM_COPY_OVERLAP" );
case -9: return QStringLiteral( "CL_IMAGE_FORMAT_MISMATCH" );
case -10: return QStringLiteral( "CL_IMAGE_FORMAT_NOT_SUPPORTED" );
case -12: return QStringLiteral( "CL_MAP_FAILURE" );
case -13: return QStringLiteral( "CL_MISALIGNED_SUB_BUFFER_OFFSET" );
case -14: return QStringLiteral( "CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST" );
case -15: return QStringLiteral( "CL_COMPILE_PROGRAM_FAILURE" );
case -16: return QStringLiteral( "CL_LINKER_NOT_AVAILABLE" );
case -17: return QStringLiteral( "CL_LINK_PROGRAM_FAILURE" );
case -18: return QStringLiteral( "CL_DEVICE_PARTITION_FAILED" );
case -19: return QStringLiteral( "CL_KERNEL_ARG_INFO_NOT_AVAILABLE" );
case -30: return QStringLiteral( "CL_INVALID_VALUE" );
case -31: return QStringLiteral( "CL_INVALID_DEVICE_TYPE" );
case -32: return QStringLiteral( "CL_INVALID_PLATFORM" );
case -33: return QStringLiteral( "CL_INVALID_DEVICE" );
case -34: return QStringLiteral( "CL_INVALID_CONTEXT" );
case -35: return QStringLiteral( "CL_INVALID_QUEUE_PROPERTIES" );
case -36: return QStringLiteral( "CL_INVALID_COMMAND_QUEUE" );
case -37: return QStringLiteral( "CL_INVALID_HOST_PTR" );
case -38: return QStringLiteral( "CL_INVALID_MEM_OBJECT" );
case -39: return QStringLiteral( "CL_INVALID_IMAGE_FORMAT_DESCRIPTOR" );
case -40: return QStringLiteral( "CL_INVALID_IMAGE_SIZE" );
case -41: return QStringLiteral( "CL_INVALID_SAMPLER" );
case -42: return QStringLiteral( "CL_INVALID_BINARY" );
case -43: return QStringLiteral( "CL_INVALID_BUILD_OPTIONS" );
case -44: return QStringLiteral( "CL_INVALID_PROGRAM" );
case -45: return QStringLiteral( "CL_INVALID_PROGRAM_EXECUTABLE" );
case -46: return QStringLiteral( "CL_INVALID_KERNEL_NAME" );
case -47: return QStringLiteral( "CL_INVALID_KERNEL_DEFINITION" );
case -48: return QStringLiteral( "CL_INVALID_KERNEL" );
case -49: return QStringLiteral( "CL_INVALID_ARG_INDEX" );
case -50: return QStringLiteral( "CL_INVALID_ARG_VALUE" );
case -51: return QStringLiteral( "CL_INVALID_ARG_SIZE" );
case -52: return QStringLiteral( "CL_INVALID_KERNEL_ARGS" );
case -53: return QStringLiteral( "CL_INVALID_WORK_DIMENSION" );
case -54: return QStringLiteral( "CL_INVALID_WORK_GROUP_SIZE" );
case -55: return QStringLiteral( "CL_INVALID_WORK_ITEM_SIZE" );
case -56: return QStringLiteral( "CL_INVALID_GLOBAL_OFFSET" );
case -57: return QStringLiteral( "CL_INVALID_EVENT_WAIT_LIST" );
case -58: return QStringLiteral( "CL_INVALID_EVENT" );
case -59: return QStringLiteral( "CL_INVALID_OPERATION" );
case -60: return QStringLiteral( "CL_INVALID_GL_OBJECT" );
case -61: return QStringLiteral( "CL_INVALID_BUFFER_SIZE" );
case -62: return QStringLiteral( "CL_INVALID_MIP_LEVEL" );
case -63: return QStringLiteral( "CL_INVALID_GLOBAL_WORK_SIZE" );
case -64: return QStringLiteral( "CL_INVALID_PROPERTY" );
case -65: return QStringLiteral( "CL_INVALID_IMAGE_DESCRIPTOR" );
case -66: return QStringLiteral( "CL_INVALID_COMPILER_OPTIONS" );
case -67: return QStringLiteral( "CL_INVALID_LINKER_OPTIONS" );
case -68: return QStringLiteral( "CL_INVALID_DEVICE_PARTITION_COUNT" );
case -69: return QStringLiteral( "CL_INVALID_PIPE_SIZE" );
case -70: return QStringLiteral( "CL_INVALID_DEVICE_QUEUE" );
case -71: return QStringLiteral( "CL_INVALID_SPEC_ID" );
case -72: return QStringLiteral( "CL_MAX_SIZE_RESTRICTION_EXCEEDED" );
case -1002: return QStringLiteral( "CL_INVALID_D3D10_DEVICE_KHR" );
case -1003: return QStringLiteral( "CL_INVALID_D3D10_RESOURCE_KHR" );
case -1004: return QStringLiteral( "CL_D3D10_RESOURCE_ALREADY_ACQUIRED_KHR" );
case -1005: return QStringLiteral( "CL_D3D10_RESOURCE_NOT_ACQUIRED_KHR" );
case -1006: return QStringLiteral( "CL_INVALID_D3D11_DEVICE_KHR" );
case -1007: return QStringLiteral( "CL_INVALID_D3D11_RESOURCE_KHR" );
case -1008: return QStringLiteral( "CL_D3D11_RESOURCE_ALREADY_ACQUIRED_KHR" );
case -1009: return QStringLiteral( "CL_D3D11_RESOURCE_NOT_ACQUIRED_KHR" );
case -1010: return QStringLiteral( "CL_INVALID_DX9_MEDIA_ADAPTER_KHR" );
case -1011: return QStringLiteral( "CL_INVALID_DX9_MEDIA_SURFACE_KHR" );
case -1012: return QStringLiteral( "CL_DX9_MEDIA_SURFACE_ALREADY_ACQUIRED_KHR" );
case -1013: return QStringLiteral( "CL_DX9_MEDIA_SURFACE_NOT_ACQUIRED_KHR" );
case -1093: return QStringLiteral( "CL_INVALID_EGL_OBJECT_KHR" );
case -1092: return QStringLiteral( "CL_EGL_RESOURCE_NOT_ACQUIRED_KHR" );
case -1001: return QStringLiteral( "CL_PLATFORM_NOT_FOUND_KHR" );
case -1057: return QStringLiteral( "CL_DEVICE_PARTITION_FAILED_EXT" );
case -1058: return QStringLiteral( "CL_INVALID_PARTITION_COUNT_EXT" );
case -1059: return QStringLiteral( "CL_INVALID_PARTITION_NAME_EXT" );
case -1094: return QStringLiteral( "CL_INVALID_ACCELERATOR_INTEL" );
case -1095: return QStringLiteral( "CL_INVALID_ACCELERATOR_TYPE_INTEL" );
case -1096: return QStringLiteral( "CL_INVALID_ACCELERATOR_DESCRIPTOR_INTEL" );
case -1097: return QStringLiteral( "CL_ACCELERATOR_TYPE_NOT_SUPPORTED_INTEL" );
case -1000: return QStringLiteral( "CL_INVALID_GL_SHAREGROUP_REFERENCE_KHR" );
case -1098: return QStringLiteral( "CL_INVALID_VA_API_MEDIA_ADAPTER_INTEL" );
case -1099: return QStringLiteral( "CL_INVALID_VA_API_MEDIA_SURFACE_INTEL" );
case -1100: return QStringLiteral( "CL_VA_API_MEDIA_SURFACE_ALREADY_ACQUIRED_INTEL" );
case -1101: return QStringLiteral( "CL_VA_API_MEDIA_SURFACE_NOT_ACQUIRED_INTEL" );
default: return QStringLiteral( "CL_UNKNOWN_ERROR" );
}
}
1 change: 1 addition & 0 deletions src/core/qgsopenclutils.h
View file
Expand Up @@ -44,6 +44,7 @@ class CORE_EXPORT QgsOpenClUtils
static QString buildLog( cl::BuildError &e );
static QString sourceFromPath( const QString &path );
static QLatin1String LOGMESSAGE_TAG;
static QString errorText( const int errorCode );

private:
QgsOpenClUtils();
Expand Down

0 comments on commit a7ef072

Please sign in to comment.