Skip to content

Commit

Permalink
Browse files Browse the repository at this point in the history
Merge pull request #9070 from elpaso/bugfix-21121-opencl-16bit-raster
Fix hillshade renderer with data type != Float32 [opencl]

Cherry-picked from master 109fcc0
  • Loading branch information
elpaso authored and nyalldawson committed Feb 6, 2019
1 parent 18c0c4d commit b292edd
Show file tree
Hide file tree
Showing 2 changed files with 44 additions and 11 deletions.
9 changes: 5 additions & 4 deletions resources/opencl_programs/hillshade_renderer.cl
@@ -1,3 +1,4 @@
// Note: "float *scanLine" may be replaced by code with actual input data type
__kernel void processNineCellWindow( __global float *scanLine1,
__global float *scanLine2,
__global float *scanLine3,
Expand Down Expand Up @@ -39,8 +40,8 @@ __kernel void processNineCellWindow( __global float *scanLine1,
if ( x32 == rasterParams[0] ) x32 = x22;
if ( x33 == rasterParams[0] ) x33 = x22;

float derX = ( ( x13 + x23 + x23 + x33 ) - ( x11 + x21 + x21 + x31 ) ) / ( 8 * rasterParams[3] );
float derY = ( ( x31 + x32 + x32 + x33 ) - ( x11 + x12 + x12 + x13 ) ) / ( 8 * -rasterParams[4]);
float derX = ( ( x13 + x23 + x23 + x33 ) - ( x11 + x21 + x21 + x31 ) ) / ( 8.0f * rasterParams[3] );
float derY = ( ( x31 + x32 + x32 + x33 ) - ( x11 + x12 + x12 + x13 ) ) / ( 8.0f * -rasterParams[4]);

if ( derX == rasterParams[0] ||
derX == rasterParams[0] )
Expand Down Expand Up @@ -90,7 +91,7 @@ __kernel void processNineCellWindow( __global float *scanLine1,
weight_270 * val270_mul_127 +
weight_315 * val315_mul_127 +
weight_360 * val360_mul_127 ) / xx_plus_yy ) /
( 1 + rasterParams[8] * xx_plus_yy );
( 1.0f + rasterParams[8] * xx_plus_yy );
res = clamp( 1.0f + cang_mul_127, 0.0f, 255.0f );
}
}
Expand All @@ -99,7 +100,7 @@ __kernel void processNineCellWindow( __global float *scanLine1,
res = ( rasterParams[9] -
( derY * rasterParams[6] -
derX * rasterParams[7] )) /
sqrt( 1 + rasterParams[8] *
sqrt( 1.0f + rasterParams[8] *
( derX * derX + derY * derY ) );
res = res <= 0.0f ? 1.0f : 1.0f + res;
}
Expand Down
46 changes: 39 additions & 7 deletions src/core/raster/qgshillshaderenderer.cpp
Expand Up @@ -30,6 +30,7 @@
#include <chrono>
#include "qgssettings.h"
#endif
#include "qgsexception.h"
#include "qgsopenclutils.h"
#endif

Expand Down Expand Up @@ -165,7 +166,8 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
bool useOpenCL( QgsOpenClUtils::enabled()
&& QgsOpenClUtils::available()
&& ( ! mRasterTransparency || mRasterTransparency->isEmpty() )
&& mAlphaBand <= 0 );
&& mAlphaBand <= 0
&& inputBlock->dataTypeSize() <= 4 );
// Check for sources
QString source;
if ( useOpenCL )
Expand All @@ -190,6 +192,36 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
std::size_t inputDataTypeSize = inputBlock->dataTypeSize();
std::size_t outputDataTypeSize = outputBlock->dataTypeSize();
// Buffer scanline, 1px height, 2px wider
QString typeName;
switch ( inputBlock->dataType() )
{
case Qgis::DataType::Byte:
typeName = QStringLiteral( "unsigned char" );
break;
case Qgis::DataType::UInt16:
typeName = QStringLiteral( "unsigned int" );
break;
case Qgis::DataType::Int16:
typeName = QStringLiteral( "short" );
break;
case Qgis::DataType::UInt32:
typeName = QStringLiteral( "unsigned int" );
break;
case Qgis::DataType::Int32:
typeName = QStringLiteral( "int" );
break;
case Qgis::DataType::Float32:
typeName = QStringLiteral( "float" );
break;
default:
throw QgsException( QStringLiteral( "Unsupported data type for OpenCL processing." ) );
}

if ( inputBlock->dataType() != Qgis::DataType::Float32 )
{
source.replace( QStringLiteral( "__global float *scanLine" ), QStringLiteral( "__global %1 *scanLine" ).arg( typeName ) );
}

// Data type for input is Float32 (4 bytes)
std::size_t scanLineWidth( inputBlock->width() + 2 );
std::size_t inputSize( inputDataTypeSize * inputBlock->width() );
Expand Down Expand Up @@ -236,7 +268,6 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
// Whether use multidirectional
rasterParams.push_back( static_cast<float>( mMultiDirectional ) ); // 17


cl::Buffer rasterParamsBuffer( queue, rasterParams.begin(), rasterParams.end(), true, false, nullptr );
cl::Buffer scanLine1Buffer( ctx, CL_MEM_READ_ONLY, bufferSize, nullptr, nullptr );
cl::Buffer scanLine2Buffer( ctx, CL_MEM_READ_ONLY, bufferSize, nullptr, nullptr );
Expand All @@ -245,13 +276,14 @@ QgsRasterBlock *QgsHillshadeRenderer::block( int bandNo, const QgsRectangle &ext
// Note that result buffer is an image
cl::Buffer resultLineBuffer( ctx, CL_MEM_WRITE_ONLY, outputDataTypeSize * width, nullptr, nullptr );

static cl::Program program;
static std::once_flag programBuilt;
std::call_once( programBuilt, [ = ]()
static std::map<Qgis::DataType, cl::Program> programCache;
cl::Program program = programCache[inputBlock->dataType()];
if ( ! program.get() )
{
// Create a program from the kernel source
program = QgsOpenClUtils::buildProgram( source, QgsOpenClUtils::ExceptionBehavior::Throw );
} );
programCache[inputBlock->dataType()] = QgsOpenClUtils::buildProgram( source, QgsOpenClUtils::ExceptionBehavior::Throw );
program = programCache[inputBlock->dataType()];
}

// Disable program cache when developing and testing cl program
// program = QgsOpenClUtils::buildProgram( ctx, source, QgsOpenClUtils::ExceptionBehavior::Throw );
Expand Down

0 comments on commit b292edd

Please sign in to comment.