From 08fc641d32ad9c8fe40fd9274493bbd736ba9c6a Mon Sep 17 00:00:00 2001 From: Alessandro Pasotti Date: Wed, 18 Apr 2018 13:47:02 +0200 Subject: [PATCH] CPLAllocator smart wrapper --- src/analysis/raster/qgsninecellfilter.cpp | 195 ++++++++++------------ src/analysis/raster/slope.cl | 1 - src/core/qgsopenclutils.cpp | 1 + src/core/qgsopenclutils.h | 57 +++++++ 4 files changed, 150 insertions(+), 104 deletions(-) diff --git a/src/analysis/raster/qgsninecellfilter.cpp b/src/analysis/raster/qgsninecellfilter.cpp index 00a694880452..54096f22dad9 100644 --- a/src/analysis/raster/qgsninecellfilter.cpp +++ b/src/analysis/raster/qgsninecellfilter.cpp @@ -221,11 +221,14 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee } //keep only three scanlines in memory at a time, make room for initial and final nodata - float *scanLine1 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) ); - float *scanLine2 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) ); - float *scanLine3 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) ); + QgsOpenClUtils::CPLAllocator scanLine1( xSize + 2 ); + QgsOpenClUtils::CPLAllocator scanLine2( xSize + 2 ); + QgsOpenClUtils::CPLAllocator scanLine3( xSize + 2 ); + //float *scanLine2 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) ); + //float *scanLine3 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) ); - float *resultLine = ( float * ) CPLMalloc( sizeof( float ) * xSize ); + //float *resultLine = ( float * ) CPLMalloc( sizeof( float ) * xSize ); + QgsOpenClUtils::CPLAllocator resultLine( xSize ); cl_int errorCode = 0; @@ -240,123 +243,109 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee addExtraRasterParams( rasterParams ); - try - { - 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 ) + 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 ) + { + if ( feedback && feedback->isCanceled() ) { - if ( feedback && feedback->isCanceled() ) - { - break; - } + break; + } - if ( feedback ) - { - feedback->setProgress( 100.0 * static_cast< double >( i ) / ySize ); - } + if ( feedback ) + { + feedback->setProgress( 100.0 * static_cast< double >( i ) / ySize ); + } - if ( i == 0 ) + 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 ) { - //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" ); - } + scanLine1[a] = mInputNodataValue; } - else + // Read scanline2 + if ( GDALRasterIO( rasterBand, GF_Read, 0, 0, xSize, 1, &scanLine2[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None ) { - //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 ) ); + QgsDebugMsg( "Raster IO Error" ); } + } + else + { + //normally fetch only scanLine3 and release scanline 1 if we move forward one row + //scanLine1 = scanLine2; + //scanLine2 = scanLine3; + //scanLine3 = ( float * ) CPLMalloc( sizeof( float ) * ( xSize + 2 ) ); + scanLine1.reset( scanLine2.release() ); + scanLine2.reset( scanLine3.release() ); + scanLine3.reset( 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 ) - { - scanLine3[a] = mInputNodataValue; - } - } - else + // Read scanline 3 + if ( i == ySize - 1 ) //fill the row below the bottom with nodata values + { + for ( int a = 0; a < xSize + 2; ++a ) { - if ( GDALRasterIO( rasterBand, GF_Read, 0, i + 1, xSize, 1, &scanLine3[1], xSize, 1, GDT_Float32, 0, 0 ) != CE_None ) - { - QgsDebugMsg( "Raster IO Error" ); - } + scanLine3[a] = mInputNodataValue; } - // 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 ) + } + 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.get() ); + errorCode = cl::enqueueWriteBuffer( scanLine2Buffer, CL_TRUE, 0, + sizeof( float ) * ( xSize + 2 ), scanLine2.get() ); + errorCode = cl::enqueueWriteBuffer( scanLine3Buffer, CL_TRUE, 0, + sizeof( float ) * ( xSize + 2 ), scanLine3.get() ); + + kernel( cl::EnqueueArgs( + cl::NDRange( xSize ) + ), + scanLine1Buffer, + scanLine2Buffer, + scanLine3Buffer, + resultLineBuffer, + rasterParamsBuffer + ); + + cl::enqueueReadBuffer( resultLineBuffer, CL_TRUE, 0, xSize * sizeof( float ), resultLine.get() ); + + if ( GDALRasterIO( outputRasterBand, GF_Write, 0, i, xSize, 1, resultLine.get(), 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 ); - CPLFree( scanLine1 ); - CPLFree( scanLine2 ); - CPLFree( scanLine3 ); if ( feedback && feedback->isCanceled() ) { diff --git a/src/analysis/raster/slope.cl b/src/analysis/raster/slope.cl index 35e58474da9c..b413ab6ff047 100644 --- a/src/analysis/raster/slope.cl +++ b/src/analysis/raster/slope.cl @@ -68,7 +68,6 @@ float calcFirstDer( float x11, float x21, float x31, float x12, float x22, float return sum / ( weight * mCellSize ) * mZFactor; } - __kernel void processNineCellWindow( __global float *scanLine1, __global float *scanLine2, __global float *scanLine3, diff --git a/src/core/qgsopenclutils.cpp b/src/core/qgsopenclutils.cpp index 5993fa6bc974..6464cfbb6201 100644 --- a/src/core/qgsopenclutils.cpp +++ b/src/core/qgsopenclutils.cpp @@ -18,6 +18,7 @@ #include "qgsmessagelog.h" #include "qgslogger.h" + #include #include #include diff --git a/src/core/qgsopenclutils.h b/src/core/qgsopenclutils.h index 5828429d7d41..12776bfcaffb 100644 --- a/src/core/qgsopenclutils.h +++ b/src/core/qgsopenclutils.h @@ -27,6 +27,7 @@ #include "qgis_core.h" #include "qgis.h" +#include "cpl_conv.h" /** * \ingroup core @@ -46,6 +47,62 @@ class CORE_EXPORT QgsOpenClUtils static QLatin1String LOGMESSAGE_TAG; static QString errorText( const int errorCode ); + /** + * Tiny smart-pointer wrapper around CPLMalloc and CPLFree: this is needed because + * OpenCL C++ API may throw exceptions + */ + template + struct CPLAllocator + { + + public: + + explicit CPLAllocator( unsigned long size ): mMem( ( T * )CPLMalloc( sizeof( T ) * size ) ) { } + + ~CPLAllocator() + { + CPLFree( ( void * )mMem ); + } + + void reset( T *newData ) + { + if ( mMem ) + CPLFree( ( void * )mMem ); + mMem = newData; + } + + void reset( unsigned long size ) + { + reset( ( T * )CPLMalloc( sizeof( T ) *size ) ); + } + + T &operator* () + { + return &mMem[0]; + } + + T *release() + { + T *tmpMem = mMem; + mMem = nullptr; + return tmpMem; + } + + T &operator[]( const int index ) + { + return mMem[index]; + } + + T *get() + { + return mMem; + } + + private: + + T *mMem = nullptr; + }; + private: QgsOpenClUtils(); static void init();