Try to avoid crash on intel haswell

Author: elpaso

src/analysis/raster/qgsninecellfilter.cpp

@@ -57,9 +57,29 @@ int QgsNineCellFilter::processRaster( QgsFeedback *feedback )
                      .arg( openClProgramBaseName( ) ), QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Info );
         return processRasterGPU( source, feedback );
       }
+<<<<<<< 583c7ae28727cddaff53706903980733bd4b8979
       catch ( cl::Error &e )
       {
         QString err = QObject::tr( "Error running OpenCL program: %1 - %2" ).arg( e.what( ) ).arg( QgsOpenClUtils::errorText( e.err( ) ) );
+=======
+      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!" );
+        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 )
+      {
+        QString err = QObject::tr( "Error %1 running OpenCL program in %2" )
+                      .arg( QgsOpenClUtils::errorText( e.err() ), QString::fromStdString( e.what() ) );
+>>>>>>> Try to avoid crash on intel haswell
         QgsMessageLog::logMessage( err, QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Critical );
         throw QgsProcessingException( err );
       }
@@ -236,6 +256,7 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
   // used to pass additional args to opencl program
   addExtraRasterParams( rasterParams );
 
+<<<<<<< 583c7ae28727cddaff53706903980733bd4b8979
   std::size_t bufferSize( sizeof( float ) * ( xSize + 2 ) );
   std::size_t inputSize( sizeof( float ) * ( xSize ) );
 
@@ -263,17 +284,40 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
 
   // 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
+>>>>>>> Try to avoid crash on intel haswell
   {
-    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;
+      }
 
+<<<<<<< 583c7ae28727cddaff53706903980733bd4b8979
     if ( i == 0 )
     {
       // Fill scanline 1 with (input) nodata for the values above the first row and
@@ -286,9 +330,35 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
 
       // Read scanline2: first real raster row
       if ( GDALRasterIO( rasterBand, GF_Read, 0, i, xSize, 1, &scanLine[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None )
+=======
+      if ( feedback )
       {
-        QgsDebugMsg( "Raster IO Error" );
+        feedback->setProgress( 100.0 * static_cast< double >( i ) / ySize );
       }
+
+      if ( i == 0 )
+>>>>>>> Try to avoid crash on intel haswell
+      {
+        //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 ) );
+      }
+<<<<<<< 583c7ae28727cddaff53706903980733bd4b8979
       queue.enqueueWriteBuffer( scanLine2Buffer, CL_TRUE, 0, bufferSize, scanLine.get() );
 
       // Read scanline3: second real raster row
@@ -303,6 +373,49 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
       // Normally fetch only scanLine3 and move forward one row
       // Read scanline 3, fill the last row with nodata values if it's the last iteration
       if ( i == ySize - 1 ) //fill the row below the bottom with nodata values
+=======
+
+      // Read scanline 3
+      if ( i == ySize - 1 ) //fill the row below the bottom with nodata values
+      {
+        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 )
+        {
+          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 )
+>>>>>>> Try to avoid crash on intel haswell
       {
         for ( int a = 0; a < xSize + 2; ++a )
         {
@@ -319,6 +432,7 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
         }
         queue.enqueueWriteBuffer( *scanLineBuffer[rowIndex[2]], CL_TRUE, 0, bufferSize, scanLine.get() ); // row 0
       }
+<<<<<<< 583c7ae28727cddaff53706903980733bd4b8979
     }
 
     kernel( cl::EnqueueArgs(
@@ -339,6 +453,18 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
       QgsDebugMsg( "Raster IO Error" );
     }
     std::rotate( rowIndex.begin(), rowIndex.begin() + 1, rowIndex.end() );
+=======
+
+    }
+  }
+  catch ( cl::Error &e )
+  {
+    CPLFree( resultLine );
+    CPLFree( scanLine1 );
+    CPLFree( scanLine2 );
+    CPLFree( scanLine3 );
+    throw e;
+>>>>>>> Try to avoid crash on intel haswell
   }
 
   if ( feedback && feedback->isCanceled() )

src/analysis/raster/slope.cl

@@ -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 )
 {
@@ -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))