Skip to content

Commit

Permalink
Merge pull request #8666 from elpaso/opencl-rastercalc-float-cast
Browse files Browse the repository at this point in the history 
opencl rastercalc fix int input rasters and cast to float
  • Loading branch information
@elpaso
elpaso committed Dec 14, 2018
2 parents 5b7f808 + b71a668 commit e58abfd
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 e58abfd

Please sign in to comment.