[opencl] Test with image2d

qgis · Aug 8, 2018 · a1a65bb · a1a65bb
1 parent 12fa896
commit a1a65bb
Show file tree

Hide file tree

Showing 6 changed files with 42 additions and 79 deletions.
diff --git a/resources/opencl_programs/aspect.cl b/resources/opencl_programs/aspect.cl
@@ -4,19 +4,13 @@ __kernel void processNineCellWindow( __global float *scanLine1,
                                      __global float *scanLine2,
                                      __global float *scanLine3,
                                      __global float *resultLine,
-<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
                                      __global float *rasterParams // mInputNodataValue, mOutputNodataValue, mZFactor, mCellSizeX, mCellSizeY
                                      )
 {
-=======
-                                     __global float *rasterParams
-                       ) {
->>>>>>> Use OpenCL command queue
 
   // Get the index of the current element
   const int i = get_global_id(0);
 
-<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
   if ( scanLine2[i+1] == rasterParams[0] )
   {
      resultLine[i] = rasterParams[1];
@@ -42,38 +36,5 @@ __kernel void processNineCellWindow( __global float *scanLine1,
     {
       resultLine[i] = 180.0f + atan2pi( derX, derY ) * 180.0f;
     }
-=======
-  // Do the operation
-  //return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX))
-  float derX = calcFirstDer(   scanLine1[i],   scanLine2[i],   scanLine3[i],
-                               scanLine1[i+1], scanLine2[i+1], scanLine3[i+1],
-                               scanLine1[i+2], scanLine2[i+2], scanLine3[i+2],
-                               rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[3]
-                             );
-  //return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * mCellSizeY));
-  float derY = calcFirstDer(   scanLine1[i+2], scanLine1[i+1], scanLine1[i],
-                               scanLine2[i+2], scanLine2[i+1], scanLine2[i],
-                               scanLine3[i+2], scanLine3[i+1], scanLine3[i],
-                               rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[4]
-                             );
-
-
-  if ( derX == rasterParams[1] || derY == rasterParams[1] ||
-       ( derX == 0.0f && derY == 0.0f) )
-  {
-    resultLine[i] = rasterParams[1];
-  }
-  else
-  {
-    // 180.0 / M_PI = 57.29577951308232
-    float aspect = atan2( derX, derY ) * 57.29577951308232;
-    if ( aspect < 0 )
-        resultLine[i] = 90.0f - aspect;
-    else if (aspect > 90.0f)
-        // 360 + 90 = 450
-        resultLine[i] = 450.0f - aspect;
-    else
-        resultLine[i] = 90.0 - aspect;
->>>>>>> Use OpenCL command queue
   }
 }
diff --git a/resources/opencl_programs/slope.cl b/resources/opencl_programs/slope.cl
@@ -4,19 +4,13 @@ __kernel void processNineCellWindow( __global float *scanLine1,
                                      __global float *scanLine2,
                                      __global float *scanLine3,
                                      __global float *resultLine,
-<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
                                      __global float *rasterParams // mInputNodataValue, mOutputNodataValue, mZFactor, mCellSizeX, mCellSizeY
                                     )
 {
-=======
-                                     __global float *rasterParams
-                       ) {
->>>>>>> Use OpenCL command queue
 
   // Get the index of the current element
   const int i = get_global_id(0);
 
-<<<<<<< 16a49cddaa18cb6d0b12335fe24c68cda183e1c0
   if ( scanLine2[i+1] == rasterParams[0] )
   {
      resultLine[i] = rasterParams[1];
@@ -44,31 +38,5 @@ __kernel void processNineCellWindow( __global float *scanLine1,
       res = atanpi( res );
       resultLine[i] = res * 180.0f;
     }
-=======
-  // Do the operation
-  //return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX))
-  float derX = calcFirstDer(   scanLine1[i],   scanLine2[i],   scanLine3[i],
-                               scanLine1[i+1], scanLine2[i+1], scanLine3[i+1],
-                               scanLine1[i+2], scanLine2[i+2], scanLine3[i+2],
-                               rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[3]
-                             );
-  //return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * mCellSizeY));
-  float derY = calcFirstDer(   scanLine1[i+2], scanLine1[i+1], scanLine1[i],
-                               scanLine2[i+2], scanLine2[i+1], scanLine2[i],
-                               scanLine3[i+2], scanLine3[i+1], scanLine3[i],
-                               rasterParams[0], rasterParams[1], rasterParams[2], rasterParams[4]
-                             );
-
-
-  if ( derX == rasterParams[1] || derY == rasterParams[1] )
-  {
-    resultLine[i] = rasterParams[1];
-  }
-  else
-  {
-    float res = sqrt( derX * derX + derY * derY );
-    res = atanpi( res );
-    resultLine[i] = res * 180.0;
->>>>>>> Use OpenCL command queue
   }
 }
diff --git a/src/analysis/raster/qgsninecellfilter.cpp b/src/analysis/raster/qgsninecellfilter.cpp
@@ -233,6 +233,11 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
 >>>>>>> Use OpenCL command queue
   QgsOpenClUtils::CPLAllocator<float> resultLine( xSize );
 
+<<<<<<< 12fa896554321892e88ca05407cf125a7ccf92c1
+=======
+  cl_int errorCode = 0;
+
+>>>>>>> [opencl] Test with image2d
   // Cast to float (because double just crashes on some GPUs)
   std::vector<float> rasterParams;
 
@@ -349,7 +354,7 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
     scanLine3[0] = scanLine3[xSize + 1] = mInputNodataValue;
 
     // TODO: There is room for further optimization here: instead of replacing the buffers
-    //       we could just replace just hthe new one (the top row) and switch the order
+    //       we could just replace just the new one (the top row) and switch the order
     //       of buffer arguments in the kernell call.
     errorCode = cl::enqueueWriteBuffer( scanLine1Buffer, CL_TRUE, 0,
                                         sizeof( float ) * ( xSize + 2 ), scanLine1.get() );

diff --git a/src/app/qgsoptions.cpp b/src/app/qgsoptions.cpp
@@ -1087,6 +1087,7 @@ QgsOptions::QgsOptions( QWidget *parent, Qt::WindowFlags fl, const QList<QgsOpti
   if ( QgsOpenClUtils::available( ) )
   {
     mGPUEnableCheckBox->setEnabled( true );
+<<<<<<< 12fa896554321892e88ca05407cf125a7ccf92c1
 
     for ( const auto &dev : QgsOpenClUtils::devices( ) )
     {
@@ -1100,13 +1101,38 @@ QgsOptions::QgsOptions( QWidget *parent, Qt::WindowFlags fl, const QList<QgsOpti
     connect( mOpenClDevicesCombo, qgis::overload< int >::of( &QComboBox::currentIndexChanged ), infoUpdater );
     mOpenClDevicesCombo->setCurrentIndex( mOpenClDevicesCombo->findData( QgsOpenClUtils::deviceId( QgsOpenClUtils::activeDevice() ) ) );
     infoUpdater( -1 );
+=======
+    mGPUInfoLabel->setText( QStringLiteral( "OpenCL compatible GPU found on your system:<br>"
+                                            "Name: <b>%1</b><br>"
+                                            "Vendor: <b>%2</b><br>"
+                                            "Profile: <b>%3</b><br>"
+                                            "Version: <b>%4</b><br>"
+                                            "Image support: <b>%5</b><br>"
+                                            "Max image2d width: <b>%6</b><br>"
+                                            "Max image2d height: <b>%7</b><br>"
+                                          ).arg( QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Name ),
+                                              QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Vendor ),
+                                              QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Profile ),
+                                              QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Version ),
+                                              QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::ImageSupport ),
+                                              QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Image2dMaxWidth ),
+                                              QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Image2dMaxHeight )
+                                               )
+                          );
+>>>>>>> [opencl] Test with image2d
   }
   else
   {
     mGPUEnableCheckBox->setEnabled( false );
+<<<<<<< 12fa896554321892e88ca05407cf125a7ccf92c1
     mGPUInfoTextBrowser->setText( tr( "An OpenCL compatible device was not found on your system.<br>"
                                       "You may need to install additional libraries in order to enable OpenCL.<br>"
                                       "Please check your logs for further details." ) );
+=======
+    mGPUInfoLabel->setText( QStringLiteral( "OpenCL compatible GPU was not found on your system.<br>"
+                                            "You may need to install additional libraries in order to enable OpenCL.<br>"
+                                            "Please check your logs for further details." ) );
+>>>>>>> [opencl] Test with image2d
   }
 
 
@@ -1662,8 +1688,12 @@ void QgsOptions::saveOptions()
 #ifdef HAVE_OPENCL
   // OpenCL settings
   QgsOpenClUtils::setEnabled( mGPUEnableCheckBox->isChecked() );
+<<<<<<< 12fa896554321892e88ca05407cf125a7ccf92c1
   QString preferredDevice( mOpenClDevicesCombo->currentData().toString() );
   QgsOpenClUtils::storePreferredDevice( preferredDevice );
+=======
+
+>>>>>>> [opencl] Test with image2d
 #endif
 
   // Gdal skip driver list

diff --git a/src/core/raster/qgscolorrampshader.cpp b/src/core/raster/qgscolorrampshader.cpp
@@ -359,6 +359,9 @@ bool QgsColorRampShader::shade( double value, int *returnRedValue, int *returnGr
     mLUTInitialized = true;
   }
 
+  if ( mLUT.empty() )
+    return false;
+
   // overflow indicates that value > maximum value + DOUBLE_DIFF_THRESHOLD
   // that way idx can point to the last valid item
   bool overflow = false;
@@ -377,10 +380,6 @@ bool QgsColorRampShader::shade( double value, int *returnRedValue, int *returnGr
       overflow = true;
     }
   }
-  else if ( lutIndex < 0 )
-  {
-    return false;
-  }
   else
   {
     // get initial value from LUT

diff --git a/src/core/raster/qgsrasterblock.h b/src/core/raster/qgsrasterblock.h
@@ -238,7 +238,7 @@ class CORE_EXPORT QgsRasterBlock
      *  \param row row index
      *  \param column column index
      *  \returns true if value is no data */
-    bool isNoData( int row, int column )
+    bool isNoData( int row, int column ) const
     {
       return isNoData( static_cast< qgssize >( row ) * mWidth + column );
     }
@@ -248,7 +248,7 @@ class CORE_EXPORT QgsRasterBlock
      *  \param row row index
      *  \param column column index
      *  \returns true if value is no data */
-    bool isNoData( qgssize row, qgssize column )
+    bool isNoData( qgssize row, qgssize column ) const
     {
       return isNoData( row * static_cast< qgssize >( mWidth ) + column );
     }
@@ -257,7 +257,7 @@ class CORE_EXPORT QgsRasterBlock
      * \brief Check if value at position is no data
      *  \param index data matrix index (long type in Python)
      *  \returns true if value is no data */
-    bool isNoData( qgssize index )
+    bool isNoData( qgssize index ) const
     {
       if ( !mHasNoDataValue && !mNoDataBitmap )
         return false;