sharpetronics/QGIS
1
0
Fork 0
mirror of https://github.com/qgis/QGIS.git synced 2025-04-14 00:07:35 -04:00
Code Issues Packages Projects Releases Wiki Activity

Use OpenCL command queue

Browse Source
This commit is contained in:
Alessandro Pasotti 2018-04-19 11:42:05 +02:00
parent 215bfd41d3
commit 3054da0c00
11 changed files with 272 additions and 255 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

3
resources/CMakeLists.txt
View File

@ -11,6 +11,9 @@ INSTALL(DIRECTORY themes DESTINATION ${QGIS_DATA_DIR}/resources)
INSTALL(DIRECTORY data DESTINATION ${QGIS_DATA_DIR}/resources)
INSTALL(DIRECTORY metadata-ISO DESTINATION ${QGIS_DATA_DIR}/resources)
INSTALL(DIRECTORY palettes DESTINATION ${QGIS_DATA_DIR}/resources)
IF (HAVE_OPENCL)
INSTALL(DIRECTORY opencl_programs DESTINATION ${QGIS_DATA_DIR}/resources)
ENDIF (HAVE_OPENCL)
IF (WITH_SERVER)
INSTALL(DIRECTORY server DESTINATION ${QGIS_DATA_DIR}/resources)

45
resources/opencl_programs/aspect.cl Normal file
View File

@ -0,0 +1,45 @@
#include "calcfirstder.cl"
__kernel void processNineCellWindow( __global float *scanLine1,
__global float *scanLine2,
__global float *scanLine3,
__global float *resultLine,
__global float *rasterParams
) {
// Get the index of the current element
const int i = get_global_id(0);
// 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;
}
}

71
resources/opencl_programs/calcfirstder.cl Normal file
View File

@ -0,0 +1,71 @@
// Calculate the first derivative from a 3x3 cell matrix
float calcFirstDer( float x11, float x21, float x31, float x12, float x22, float x32, float x13, float x23, float x33,
float inputNodataValue, float outputNodataValue, float zFactor, float mCellSize )
{
//the basic formula would be simple, but we need to test for nodata values...
//X: return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * cellSizeX));
//Y: return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * cellSizeY));
int weight = 0;
float sum = 0;
//first row
if ( x31 != inputNodataValue && x11 != inputNodataValue ) //the normal case
{
sum += ( x31 - x11 );
weight += 2;
}
else if ( x31 == inputNodataValue && x11 != inputNodataValue && x21 != inputNodataValue ) //probably 3x3 window is at the border
{
sum += ( x21 - x11 );
weight += 1;
}
else if ( x11 == inputNodataValue && x31 != inputNodataValue && x21 != inputNodataValue ) //probably 3x3 window is at the border
{
sum += ( x31 - x21 );
weight += 1;
}
//second row
if ( x32 != inputNodataValue && x12 != inputNodataValue ) //the normal case
{
sum += 2.0f * ( x32 - x12 );
weight += 4;
}
else if ( x32 == inputNodataValue && x12 != inputNodataValue && x22 != inputNodataValue )
{
sum += 2.0f * ( x22 - x12 );
weight += 2;
}
else if ( x12 == inputNodataValue && x32 != inputNodataValue && x22 != inputNodataValue )
{
sum += 2.0f * ( x32 - x22 );
weight += 2;
}
//third row
if ( x33 != inputNodataValue && x13 != inputNodataValue ) //the normal case
{
sum += ( x33 - x13 );
weight += 2;
}
else if ( x33 == inputNodataValue && x13 != inputNodataValue && x23 != inputNodataValue )
{
sum += ( x23 - x13 );
weight += 1;
}
else if ( x13 == inputNodataValue && x33 != inputNodataValue && x23 != inputNodataValue )
{
sum += ( x33 - x23 );
weight += 1;
}
if ( weight == 0 )
{
return outputNodataValue;
}
return sum / ( weight * mCellSize ) * zFactor;
}

38
resources/opencl_programs/slope.cl Normal file
View File

@ -0,0 +1,38 @@
#include "calcfirstder.cl"
__kernel void processNineCellWindow( __global float *scanLine1,
__global float *scanLine2,
__global float *scanLine3,
__global float *resultLine,
__global float *rasterParams
) {
// Get the index of the current element
const int i = get_global_id(0);
// 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;
}
}

114
src/analysis/raster/aspect.cl
View File

@ -1,114 +0,0 @@
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 )
{
//the basic formula would be simple, but we need to test for nodata values...
//X: return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX));
//Y: return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * mCellSizeY));
int weight = 0;
float sum = 0;
//first row
if ( x31 != mInputNodataValue && x11 != mInputNodataValue ) //the normal case
{
sum += ( x31 - x11 );
weight += 2;
}
else if ( x31 == mInputNodataValue && x11 != mInputNodataValue && x21 != mInputNodataValue ) //probably 3x3 window is at the border
{
sum += ( x21 - x11 );
weight += 1;
}
else if ( x11 == mInputNodataValue && x31 != mInputNodataValue && x21 != mInputNodataValue ) //probably 3x3 window is at the border
{
sum += ( x31 - x21 );
weight += 1;
}
//second row
if ( x32 != mInputNodataValue && x12 != mInputNodataValue ) //the normal case
{
sum += 2.0f * ( x32 - x12 );
weight += 4;
}
else if ( x32 == mInputNodataValue && x12 != mInputNodataValue && x22 != mInputNodataValue )
{
sum += 2.0f * ( x22 - x12 );
weight += 2;
}
else if ( x12 == mInputNodataValue && x32 != mInputNodataValue && x22 != mInputNodataValue )
{
sum += 2.0f * ( x32 - x22 );
weight += 2;
}
//third row
if ( x33 != mInputNodataValue && x13 != mInputNodataValue ) //the normal case
{
sum += ( x33 - x13 );
weight += 2;
}
else if ( x33 == mInputNodataValue && x13 != mInputNodataValue && x23 != mInputNodataValue )
{
sum += ( x23 - x13 );
weight += 1;
}
else if ( x13 == mInputNodataValue && x33 != mInputNodataValue && x23 != mInputNodataValue )
{
sum += ( x33 - x23 );
weight += 1;
}
if ( weight == 0 )
{
return mOutputNodataValue;
}
return sum / ( weight * mCellSize ) * mZFactor;
}
__kernel void processNineCellWindow( __global float *scanLine1,
__global float *scanLine2,
__global float *scanLine3,
__global float *resultLine,
__global float *rasterParams
) {
// Get the index of the current element
const int i = get_global_id(0);
// 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;
}
}

33
src/analysis/raster/qgsninecellfilter.cpp
View File

@ -41,7 +41,7 @@ int QgsNineCellFilter::processRaster( QgsFeedback *feedback )
if ( QgsOpenClUtils::enabled() && QgsOpenClUtils::available() && ! openClProgramBaseName( ).isEmpty() )
{
// Load the program sources
QString source( QgsOpenClUtils::sourceFromPath( QStringLiteral( "/home/ale/dev/QGIS/src/analysis/raster/%1.cl" ).arg( openClProgramBaseName( ) ) ) );
QString source( QgsOpenClUtils::sourceFromBaseName( openClProgramBaseName( ) ) );
if ( ! source.isEmpty() )
{
try
@ -50,19 +50,6 @@ int QgsNineCellFilter::processRaster( QgsFeedback *feedback )
.arg( openClProgramBaseName( ) ), QgsOpenClUtils::LOGMESSAGE_TAG, Qgis::Info );
return processRasterGPU( source, feedback );
}
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" )
@ -220,18 +207,15 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
return 6;
}
// Prepare context
// Prepare context and queue
cl::Context ctx = QgsOpenClUtils::context();
cl::Context::setDefault( ctx );
cl::CommandQueue queue( ctx );
//keep only three scanlines in memory at a time, make room for initial and final nodata
QgsOpenClUtils::CPLAllocator<float> scanLine1( xSize + 2 );
QgsOpenClUtils::CPLAllocator<float> scanLine2( xSize + 2 );
QgsOpenClUtils::CPLAllocator<float> 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 );
QgsOpenClUtils::CPLAllocator<float> resultLine( xSize );
cl_int errorCode = 0;
@ -255,9 +239,7 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
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" );
cl::Program program( QgsOpenClUtils::buildProgram( ctx, source, QgsOpenClUtils::ExceptionBehavior::Throw ) );
// Create the OpenCL kernel
auto kernel = cl::KernelFunctor <
@ -297,9 +279,6 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
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 );
@ -325,6 +304,9 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
scanLine2[0] = scanLine2[xSize + 1] = mInputNodataValue;
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
// of buffer arguments in the kernell call.
errorCode = cl::enqueueWriteBuffer( scanLine1Buffer, CL_TRUE, 0,
sizeof( float ) * ( xSize + 2 ), scanLine1.get() );
errorCode = cl::enqueueWriteBuffer( scanLine2Buffer, CL_TRUE, 0,
@ -333,6 +315,7 @@ int QgsNineCellFilter::processRasterGPU( const QString &source, QgsFeedback *fee
sizeof( float ) * ( xSize + 2 ), scanLine3.get() );
kernel( cl::EnqueueArgs(
queue,
cl::NDRange( xSize )
),
scanLine1Buffer,

106
src/analysis/raster/slope.cl
View File

@ -1,106 +0,0 @@
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 )
{
//the basic formula would be simple, but we need to test for nodata values...
//X: return (( (x31 - x11) + 2 * (x32 - x12) + (x33 - x13) ) / (8 * mCellSizeX));
//Y: return (((x11 - x13) + 2 * (x21 - x23) + (x31 - x33)) / ( 8 * mCellSizeY));
int weight = 0;
float sum = 0;
//first row
if ( x31 != mInputNodataValue && x11 != mInputNodataValue ) //the normal case
{
sum += ( x31 - x11 );
weight += 2;
}
else if ( x31 == mInputNodataValue && x11 != mInputNodataValue && x21 != mInputNodataValue ) //probably 3x3 window is at the border
{
sum += ( x21 - x11 );
weight += 1;
}
else if ( x11 == mInputNodataValue && x31 != mInputNodataValue && x21 != mInputNodataValue ) //probably 3x3 window is at the border
{
sum += ( x31 - x21 );
weight += 1;
}
//second row
if ( x32 != mInputNodataValue && x12 != mInputNodataValue ) //the normal case
{
sum += 2.0f * ( x32 - x12 );
weight += 4;
}
else if ( x32 == mInputNodataValue && x12 != mInputNodataValue && x22 != mInputNodataValue )
{
sum += 2.0f * ( x22 - x12 );
weight += 2;
}
else if ( x12 == mInputNodataValue && x32 != mInputNodataValue && x22 != mInputNodataValue )
{
sum += 2.0f * ( x32 - x22 );
weight += 2;
}
//third row
if ( x33 != mInputNodataValue && x13 != mInputNodataValue ) //the normal case
{
sum += ( x33 - x13 );
weight += 2;
}
else if ( x33 == mInputNodataValue && x13 != mInputNodataValue && x23 != mInputNodataValue )
{
sum += ( x23 - x13 );
weight += 1;
}
else if ( x13 == mInputNodataValue && x33 != mInputNodataValue && x23 != mInputNodataValue )
{
sum += ( x33 - x23 );
weight += 1;
}
if ( weight == 0 )
{
return mOutputNodataValue;
}
return sum / ( weight * mCellSize ) * mZFactor;
}
__kernel void processNineCellWindow( __global float *scanLine1,
__global float *scanLine2,
__global float *scanLine3,
__global float *resultLine,
__global float *rasterParams
) {
// Get the index of the current element
const int i = get_global_id(0);
// 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;
}
}

11
src/app/qgisapp.cpp
View File

@ -96,6 +96,9 @@
#include "qgsgui.h"
#include "qgsnative.h"
#ifdef HAVE_OPENCL
#include "qgsopenclutils.h"
#endif
#include <QNetworkReply>
#include <QNetworkProxy>
@ -1110,6 +1113,14 @@ QgisApp::QgisApp( QSplashScreen *splash, bool restorePlugins, bool skipVersionCh
// Setup QgsNetworkAccessManager (this needs to happen after authentication, for proxy settings)
namSetup();
#ifdef HAVE_OPENCL
// Setup OpenCL source path
// TODO: cmd line switch and env var
QgsOpenClUtils::setSourcePath( QDir( QgsApplication::pkgDataPath() ).absoluteFilePath( QStringLiteral( "resources/opencl_programs" ) ) );
#endif
QgsApplication::dataItemProviderRegistry()->addProvider( new QgsQlrDataItemProvider() );
registerCustomDropHandler( new QgsQlrDropHandler() );
QgsApplication::dataItemProviderRegistry()->addProvider( new QgsQptDataItemProvider() );

28
src/core/qgsopenclutils.cpp
View File

@ -28,6 +28,7 @@ QLatin1String QgsOpenClUtils::LOGMESSAGE_TAG = QLatin1Literal( "OpenCL" );
bool QgsOpenClUtils::sAvailable = false;
cl::Platform QgsOpenClUtils::sPlatform = cl::Platform();
cl::Device QgsOpenClUtils::sDevice = cl::Device();
QString QgsOpenClUtils::sSourcePath = QString();
void QgsOpenClUtils::init()
{
@ -103,6 +104,16 @@ void QgsOpenClUtils::init()
}
}
QString QgsOpenClUtils::sourcePath()
{
return sSourcePath.isEmpty() ? QStringLiteral( "./" ) : sSourcePath;
}
void QgsOpenClUtils::setSourcePath( const QString &value )
{
sSourcePath = value;
}
bool QgsOpenClUtils::enabled()
{
@ -139,6 +150,12 @@ QString QgsOpenClUtils::sourceFromPath( const QString &path )
return source_str;
}
QString QgsOpenClUtils::sourceFromBaseName( const QString &baseName )
{
QString path = QStringLiteral( "%1/%2.cl" ).arg( sourcePath(), baseName );
return sourceFromPath( path );
}
QString QgsOpenClUtils::buildLog( cl::BuildError &e )
{
cl::BuildLogType build_logs = e.getBuildLog();
@ -263,15 +280,14 @@ cl::Program QgsOpenClUtils::buildProgram( const cl::Context &context, const QStr
try
{
program = cl::Program( context, source.toStdString( ) );
program.build( "-cl-std=CL1.1" );
// OpenCL 1.1 for compatibility with older hardware
// TODO: make this configurable
program.build( QStringLiteral( "-cl-std=CL1.1 -I%1" ).arg( sourcePath() ).toStdString().c_str() );
}
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
QString build_log( buildLog( e ) );
if ( build_log.isEmpty() )
build_log = QObject::tr( "Build logs not available!" );
QString err = QObject::tr( "Error building OpenCL program: %1" )
.arg( build_log );

18
src/core/qgsopenclutils.h
View File

@ -32,7 +32,12 @@
/**
* \ingroup core
* \class QgsOpenClUtils
* \brief The QgsOpenClUtils class is responsible for common OpenCL operations
* \brief The QgsOpenClUtils class is responsible for common OpenCL operations such as
* - enable/disable opencl
* - check opencl device availability and automatically choose the first GPU (TODO: let the user choose & override!)
* - creating contexts
* - loading program sources from standard locations
* - build programs and log errors
* \since QGIS 3.4
* \note not available in Python bindings
*/
@ -42,8 +47,8 @@ class CORE_EXPORT QgsOpenClUtils
enum ExceptionBehavior
{
Catch,
Throw
Catch, // Write errors in the message log and silently fail
Throw // Write errors in the message log and re-throw exceptions
};
static bool enabled();
@ -51,13 +56,16 @@ class CORE_EXPORT QgsOpenClUtils
static void setEnabled( bool enabled );
static QString buildLog( cl::BuildError &e );
static QString sourceFromPath( const QString &path );
static QString sourceFromBaseName( const QString &baseName );
static QLatin1String LOGMESSAGE_TAG;
static QString errorText( const int errorCode );
static cl::Program buildProgram( const cl::Context &context, const QString &source, ExceptionBehavior exceptionBehavior = Catch );
static cl::Context context();
static QString sourcePath();
static void setSourcePath( const QString &value );
/**
* Tiny smart-pointer wrapper around CPLMalloc and CPLFree: this is needed because
* Tiny smart-pointer-like wrapper around CPLMalloc and CPLFree: this is needed because
* OpenCL C++ API may throw exceptions
*/
template <typename T>
@ -112,6 +120,7 @@ class CORE_EXPORT QgsOpenClUtils
T *mMem = nullptr;
};
private:
QgsOpenClUtils();
static void init();
@ -119,6 +128,7 @@ class CORE_EXPORT QgsOpenClUtils
static cl::Device sDevice;
static cl::Platform sPlatform;
static QLatin1String SETTINGS_KEY;
static QString sSourcePath;
};

60
tests/src/core/testqgsopenclutils.cpp
View File

@ -17,6 +17,9 @@
#include "qgstest.h"
#include <QObject>
#include <QString>
#include <QTemporaryFile>
#include <qgsapplication.h>
//header for class being tested
#include <qgsopenclutils.h>
@ -29,15 +32,22 @@ class TestQgsOpenClUtils: public QObject
//void testRunMakeProgram();
private slots:
void initTestCase();// will be called before the first testfunction is executed.
void cleanupTestCase();// will be called after the last testfunction was executed.
void init();// will be called before each testfunction is executed.
void cleanup();// will be called after every testfunction.
void TestEnable();
void TestDisable();
void TestAvailable();
void testMakeRunProgram();
void testProgramSource();
void testContext();
private:
void _testMakeRunProgram();
cl::Program buildProgram( const cl::Context &context, const QString &source )
{
cl::Program program( context, source.toStdString( ) );
@ -59,6 +69,36 @@ class TestQgsOpenClUtils: public QObject
}
};
void TestQgsOpenClUtils::init()
{
}
void TestQgsOpenClUtils::cleanup()
{
}
void TestQgsOpenClUtils::initTestCase()
{
// Runs once before any tests are run
// Set up the QgsSettings environment
QCoreApplication::setOrganizationName( QStringLiteral( "QGIS" ) );
QCoreApplication::setOrganizationDomain( QStringLiteral( "qgis.org" ) );
QCoreApplication::setApplicationName( QStringLiteral( "QGIS-TEST" ) );
QgsApplication::init();
QgsApplication::initQgis();
}
void TestQgsOpenClUtils::cleanupTestCase()
{
// Runs once after all tests are run
QgsApplication::exitQgis();
}
void TestQgsOpenClUtils::TestEnable()
{
QgsOpenClUtils::setEnabled( true );
@ -78,6 +118,13 @@ void TestQgsOpenClUtils::TestAvailable()
void TestQgsOpenClUtils::testMakeRunProgram()
{
// Run twice to check for valid command queue in the second call
_testMakeRunProgram();
_testMakeRunProgram();
}
void TestQgsOpenClUtils::_testMakeRunProgram()
{
cl_int err = 0;
@ -86,6 +133,7 @@ void TestQgsOpenClUtils::testMakeRunProgram()
cl::Context ctx = QgsOpenClUtils::context();
cl::Context::setDefault( ctx );
cl::CommandQueue queue( ctx );
std::vector<float> a_vec = {1, 10, 100};
std::vector<float> b_vec = {1, 10, 100};
@ -104,6 +152,7 @@ void TestQgsOpenClUtils::testMakeRunProgram()
> ( program, "vectorAdd" );
kernel( cl::EnqueueArgs(
queue,
cl::NDRange( 3 )
),
a_buf,
@ -118,6 +167,17 @@ void TestQgsOpenClUtils::testMakeRunProgram()
}
}
void TestQgsOpenClUtils::testProgramSource()
{
QgsOpenClUtils::setSourcePath( QDir::tempPath() );
QTemporaryFile tmpFile( QDir::tempPath() + "/XXXXXX.cl" );
tmpFile.open( );
tmpFile.write( QByteArray::fromStdString( source( ) ) );
tmpFile.flush();
QString baseName = tmpFile.fileName().replace( ".cl", "" ).replace( QDir::tempPath(), "" );
QVERIFY( ! QgsOpenClUtils::sourceFromBaseName( baseName ).isEmpty() );
}
void TestQgsOpenClUtils::testContext()
{
QVERIFY( QgsOpenClUtils::context()() != nullptr );