Skip to content
Permalink
Browse files

opencl rastercalc fix int input rasters and cast to float

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 committed Dec 13, 2018
1 parent 8c07c99 commit b71a668a0892b1a9a1c877c05df3310c59f23860
@@ -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:
@@ -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;
@@ -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;
@@ -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" );
@@ -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
)
{
@@ -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 );
@@ -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();
@@ -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() );

@@ -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 ) )" ) );
}


0 comments on commit b71a668

Please sign in to comment.
You can’t perform that action at this time.