Skip to content

Commit

Permalink
opencl rastercalc fix int input rasters and cast to float
Browse files Browse the repository at this point in the history 
Cast to float all math operations because when the
input is not a float or a double opencl raises
an error regarding which override should pick.

By casting to float we are sure that the right
function will be called.

This patch also fixes the buffer sizes for short (16bit)
and int (32bit) and asserts that siexe of float is 32bit.
  • Loading branch information
@elpaso
elpaso committed Dec 13, 2018
1 parent 8c07c99 commit b71a668
Show file tree
Hide file tree
Showing 3 changed files with 55 additions and 19 deletions.
53 changes: 43 additions & 10 deletions src/analysis/raster/qgsrastercalcnode.cpp
View file
Expand Up @@ -207,6 +207,12 @@ QString QgsRasterCalcNode::toString( bool cStyle ) const
left = mLeft->toString( cStyle );
if ( mRight )
right = mRight->toString( cStyle );

auto floatCast = [ ]( const QString s ) -> QString
{
return QStringLiteral( "(float) ( %1 )" ).arg( s );
};

switch ( mType )
{
case tOperator:
Expand All @@ -227,7 +233,7 @@ QString QgsRasterCalcNode::toString( bool cStyle ) const
break;
case opPOW:
if ( cStyle )
result = QStringLiteral( "pow( %1, %2 )" ).arg( left ).arg( right );
result = QStringLiteral( "pow( %1, %2 )" ).arg( floatCast( left ) ).arg( floatCast( right ) );
else
result = QStringLiteral( "%1^%2" ).arg( left ).arg( right );
break;
Expand Down Expand Up @@ -265,31 +271,58 @@ QString QgsRasterCalcNode::toString( bool cStyle ) const
result = QStringLiteral( "%1 OR %2" ).arg( left ).arg( right );
break;
case opSQRT:
result = QStringLiteral( "sqrt( %1 )" ).arg( left );
if ( cStyle )
result = QStringLiteral( "sqrt( %1 )" ).arg( floatCast( left ) );
else
result = QStringLiteral( "sqrt( %1 )" ).arg( left );
break;
case opSIN:
result = QStringLiteral( "sin( %1 )" ).arg( left );
if ( cStyle )
result = QStringLiteral( "sin( %1 )" ).arg( floatCast( left ) );
else
result = QStringLiteral( "sin( %1 )" ).arg( left );
break;
case opCOS:
result = QStringLiteral( "cos( %1 )" ).arg( left );
if ( cStyle )
result = QStringLiteral( "cos( %1 )" ).arg( floatCast( left ) );
else
result = QStringLiteral( "cos( %1 )" ).arg( left );
break;
case opTAN:
result = QStringLiteral( "tan( %1 )" ).arg( left );
if ( cStyle )
result = QStringLiteral( "tan( %1 )" ).arg( floatCast( left ) );
else
result = QStringLiteral( "tan( %1 )" ).arg( left );
break;
case opASIN:
result = QStringLiteral( "asin( %1 )" ).arg( left );
if ( cStyle )
result = QStringLiteral( "asin( %1 )" ).arg( floatCast( left ) );
else
result = QStringLiteral( "asin( %1 )" ).arg( left );
break;
case opACOS:
result = QStringLiteral( "acos( %1 )" ).arg( left );
if ( cStyle )
result = QStringLiteral( "acos( %1 )" ).arg( floatCast( left ) );
else
result = QStringLiteral( "acos( %1 )" ).arg( left );
break;
case opATAN:
result = QStringLiteral( "atan( %1 )" ).arg( left );
if ( cStyle )
result = QStringLiteral( "atan( %1 )" ).arg( floatCast( left ) );
else
result = QStringLiteral( "atan( %1 )" ).arg( left );
break;
case opLOG:
result = QStringLiteral( "log( %1 )" ).arg( left );
if ( cStyle )
result = QStringLiteral( "log( %1 )" ).arg( floatCast( left ) );
else
result = QStringLiteral( "log( %1 )" ).arg( left );
break;
case opLOG10:
result = QStringLiteral( "log10( %1 )" ).arg( left );
if ( cStyle )
result = QStringLiteral( "log10( %1 )" ).arg( floatCast( left ) );
else
result = QStringLiteral( "log10( %1 )" ).arg( left );
break;
case opNONE:
break;
Expand Down
17 changes: 10 additions & 7 deletions src/analysis/raster/qgsrastercalculator.cpp
View file
Expand Up @@ -370,13 +370,13 @@ QgsRasterCalculator::Result QgsRasterCalculator::processCalculationGPU( std::uni
entry.typeName = QStringLiteral( "unsigned int" );
break;
case Qgis::DataType::Int16:
entry.typeName = QStringLiteral( "int" );
entry.typeName = QStringLiteral( "short" );
break;
case Qgis::DataType::UInt32:
entry.typeName = QStringLiteral( "unsigned long" );
entry.typeName = QStringLiteral( "unsigned int" );
break;
case Qgis::DataType::Int32:
entry.typeName = QStringLiteral( "long" );
entry.typeName = QStringLiteral( "int" );
break;
case Qgis::DataType::Float32:
entry.typeName = QStringLiteral( "float" );
Expand Down Expand Up @@ -422,7 +422,7 @@ QgsRasterCalculator::Result QgsRasterCalculator::processCalculationGPU( std::uni
// Inputs:
##INPUT_DESC##
// Expression: ##EXPRESSION_ORIGINAL##
__kernel void rasterCalculator( ##INPUT##,
__kernel void rasterCalculator( ##INPUT##
__global float *resultLine
)
{
Expand All @@ -439,16 +439,18 @@ QgsRasterCalculator::Result QgsRasterCalculator::processCalculationGPU( std::uni
inputDesc.append( QStringLiteral( " // %1 = %2" ).arg( ref.varName ).arg( ref.name ) );
}
programTemplate = programTemplate.replace( QStringLiteral( "##INPUT_DESC##" ), inputDesc.join( '\n' ) );
programTemplate = programTemplate.replace( QStringLiteral( "##INPUT##" ), inputArgs.join( ',' ) );
programTemplate = programTemplate.replace( QStringLiteral( "##INPUT##" ), inputArgs.length() ? ( inputArgs.join( ',' ).append( ',' ) ) : QChar( ' ' ) );
programTemplate = programTemplate.replace( QStringLiteral( "##EXPRESSION##" ), cExpression );
programTemplate = programTemplate.replace( QStringLiteral( "##EXPRESSION_ORIGINAL##" ), calcNode->toString( ) );

//qDebug() << programTemplate;
qDebug() << programTemplate;

// Create a program from the kernel source
cl::Program program( QgsOpenClUtils::buildProgram( programTemplate, QgsOpenClUtils::ExceptionBehavior::Throw ) );

// Create the buffers, output is float32 (4 bytes)
// We assume size of float = 4 because that's the size used by OpenCL and IEEE 754
Q_ASSERT( sizeof( float ) == 4 );
std::size_t resultBufferSize( 4 * static_cast<size_t>( mNumOutputColumns ) );
cl::Buffer resultLineBuffer( ctx, CL_MEM_WRITE_ONLY,
resultBufferSize, nullptr, nullptr );
Expand All @@ -461,7 +463,7 @@ QgsRasterCalculator::Result QgsRasterCalculator::processCalculationGPU( std::uni
}
kernel.setArg( static_cast<unsigned int>( inputBuffers.size() ), resultLineBuffer );

QgsOpenClUtils::CPLAllocator<float> resultLine( resultBufferSize );
QgsOpenClUtils::CPLAllocator<float> resultLine( static_cast<size_t>( mNumOutputColumns ) );

//open output dataset for writing
GDALDriverH outputDriver = openOutputDriver();
Expand Down Expand Up @@ -528,6 +530,7 @@ QgsRasterCalculator::Result QgsRasterCalculator::processCalculationGPU( std::uni
// qDebug() << "Input: " << line << i << ref.varName << " = " << block->value( 0, i );
//qDebug() << "Writing buffer " << ref.index;

Q_ASSERT( ref.bufferSize == static_cast<size_t>( block->data().size( ) ) );
queue.enqueueWriteBuffer( inputBuffers[ref.index], CL_TRUE, 0,
ref.bufferSize, block->bits() );

Expand Down
4 changes: 2 additions & 2 deletions tests/src/analysis/testqgsrastercalculator.cpp
View file
Expand Up @@ -679,9 +679,9 @@ void TestQgsRasterCalculator::toString()
QCOMPARE( _test( QStringLiteral( "\"raster@1\" + 2" ), false ), QString( "\"raster@1\" + 2" ) );
QCOMPARE( _test( QStringLiteral( "\"raster@1\" + 2" ), true ), QString( "\"raster@1\" + 2" ) );
QCOMPARE( _test( QStringLiteral( "\"raster@1\" ^ 3 + 2" ), false ), QString( "\"raster@1\"^3 + 2" ) );
QCOMPARE( _test( QStringLiteral( "\"raster@1\" ^ 3 + 2" ), true ), QString( "pow( \"raster@1\", 3 ) + 2" ) );
QCOMPARE( _test( QStringLiteral( "\"raster@1\" ^ 3 + 2" ), true ), QString( "pow( (float) ( \"raster@1\" ), (float) ( 3 ) ) + 2" ) );
QCOMPARE( _test( QStringLiteral( "atan(\"raster@1\") * cos( 3 + 2 )" ), false ), QString( "atan( \"raster@1\" ) * cos( 3 + 2 )" ) );
QCOMPARE( _test( QStringLiteral( "atan(\"raster@1\") * cos( 3 + 2 )" ), true ), QString( "atan( \"raster@1\" ) * cos( 3 + 2 )" ) );
QCOMPARE( _test( QStringLiteral( "atan(\"raster@1\") * cos( 3 + 2 )" ), true ), QString( "atan( (float) ( \"raster@1\" ) ) * cos( (float) ( 3 + 2 ) )" ) );
}


Expand Down

0 comments on commit b71a668

Please sign in to comment.