[opencl] Test with image2d

resources/opencl_programs/hillshade.cl

@@ -0,0 +1,35 @@
+#include "calcfirstder.cl"
+
+__constant sampler_t sampler =
+      CLK_NORMALIZED_COORDS_FALSE
+    | CLK_ADDRESS_CLAMP_TO_EDGE
+    | CLK_FILTER_NEAREST;
+
+__kernel void processNineCellWindow(
+                                     __read_only image2d_t inputImage,
+                                     image2d_t outputImage,
+                                     __global float *rasterParams
+                       ) {
+
+
+  // 0=width, 1=height
+  int2 currentPosition = (int2)(get_global_id(0), get_global_id(1));
+  float4 currentPixel = (ufloat4)(0.0f);
+  float4 calculatedPixel = (float4)(1.0f);
+
+  currentPixel = read_imageuf(inputImage, sampler, currentPosition);
+  calculatedPixel = currentPixel;
+  write_imageuf(outputImage, currentPosition, calculatedPixel);
+}
+
+
+const sampler_t smp = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_LINEAR;
+
+void kernel copy(__read_only image2d_t in, image2d_t out)
+{
+  int x = get_global_id(0);
+  int y = get_global_id(1);
+  int2 pos = (int2)(x, y);
+  uint4 pixel = read_imageui(in, smp, pos);
+  write_imageui(out, pos, pixel);
+}

src/analysis/raster/qgsninecellfilter.cpp

@@ -220,7 +220,7 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
 
   cl_int errorCode = 0;
 
-  // Cast to float
+  // Cast to float (because double just crashes on some GPUs)
   std::vector<float> rasterParams;
 
   rasterParams.push_back( mInputNodataValue );
@@ -305,7 +305,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() );

src/app/qgsoptions.cpp

@@ -1092,20 +1092,25 @@ QgsOptions::QgsOptions( QWidget *parent, Qt::WindowFlags fl, const QList<QgsOpti
                                             "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::Version ),
+                                              QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::ImageSupport ),
+                                              QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Image2dMaxWidth ),
+                                              QgsOpenClUtils::deviceInfo( QgsOpenClUtils::Info::Image2dMaxHeight )
+                                               )
                           );
-    connect( mGPUEnableCheckBox, &QCheckBox::toggled, this, []( bool status )
-    {
-      QgsOpenClUtils::setEnabled( status );
-    }, Qt::UniqueConnection );
   }
   else
   {
     mGPUEnableCheckBox->setEnabled( false );
-    mGPUInfoLabel->setText( QStringLiteral( "OpenCL compatible GPU was not found on your system. You may need to install additional libraries in order to enable OpenCL." ) );
+    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." ) );
   }
 
 
@@ -1658,6 +1663,12 @@ void QgsOptions::saveOptions()
   // Number settings
   mSettings->setValue( QStringLiteral( "locale/showGroupSeparator" ), cbShowGroupSeparator->isChecked( ) );
 
+#ifdef HAVE_OPENCL
+  // OpenCL settings
+  QgsOpenClUtils::setEnabled( mGPUEnableCheckBox->isChecked() );
+
+#endif
+
   // Gdal skip driver list
   if ( mLoadedGdalDriverList )
     saveGdalDriverList();

src/core/qgsopenclutils.cpp

@@ -126,6 +126,12 @@ QString QgsOpenClUtils::deviceInfo( const Info infoType )
         return QString::fromStdString( sDevice.getInfo<CL_DEVICE_PROFILE>() );
       case Info::Version:
         return QString::fromStdString( sDevice.getInfo<CL_DEVICE_VERSION>() );
+      case Info::ImageSupport:
+        return sDevice.getInfo<CL_DEVICE_IMAGE_SUPPORT>() ? QStringLiteral( "True" ) : QStringLiteral( "False" );
+      case Info::Image2dMaxHeight:
+        return QString::number( sDevice.getInfo<CL_DEVICE_IMAGE2D_MAX_HEIGHT>() );
+      case Info::Image2dMaxWidth:
+        return QString::number( sDevice.getInfo<CL_DEVICE_IMAGE2D_MAX_WIDTH>() );
       case Info::Name:
       default:
         return QString::fromStdString( sDevice.getInfo<CL_DEVICE_NAME>() );
@@ -151,6 +157,8 @@ void QgsOpenClUtils::setEnabled( bool enabled )
   QgsSettings().setValue( SETTINGS_KEY, enabled, QgsSettings::Section::Core );
 }
 
+
+
 QString QgsOpenClUtils::sourceFromPath( const QString &path )
 {
   // TODO: check for compatibility with current platform ( cl_khr_fp64 )

src/core/qgsopenclutils.h

@@ -55,7 +55,10 @@ class CORE_EXPORT QgsOpenClUtils
       Name = CL_DEVICE_NAME,
       Vendor = CL_DEVICE_VENDOR,
       Version = CL_DEVICE_VERSION,
-      Profile = CL_DEVICE_PROFILE
+      Profile = CL_DEVICE_PROFILE,
+      ImageSupport = CL_DEVICE_IMAGE_SUPPORT,
+      Image2dMaxWidth = CL_DEVICE_IMAGE2D_MAX_WIDTH,
+      Image2dMaxHeight = CL_DEVICE_IMAGE2D_MAX_HEIGHT
     };
 
     static bool enabled();
@@ -67,6 +70,14 @@ class CORE_EXPORT QgsOpenClUtils
     static QLatin1String LOGMESSAGE_TAG;
     static QString errorText( const int errorCode );
     static cl::Program buildProgram( const cl::Context &context, const QString &source, ExceptionBehavior exceptionBehavior = Catch );
+
+    /**
+     * Context factory
+     *
+     * \return a new context for the default device or an invalid context if
+     *         no device were identified or OpenCL support is not available
+     *         and enabled
+     */
     static cl::Context context();
     static QString sourcePath();
     static void setSourcePath( const QString &value );